// SPDX-License-Identifier: GPL-3.0-or-later
#include <chrono>
#include "lifter/bridge.h"
#include "lifter/gcn.h"
#include "lifter/lift.h"
#include <algorithm>
#include <atomic>
#include <climits>
#include <cstdint>
#include <mutex>
#include <thread>
#include <array>
#include <cstddef>
#include <cstdio>
#include <cstring>
#include <filesystem>
#include <fstream>
#include <map>
#include <new>
#include <string>
#include <vector>
#ifdef _WIN32
#define WIN32_LEAN_AND_MEAN
#include <windows.h>
#else
#include <sys/wait.h>
#endif

#include "shader_recompiler/backend/bindings.h"
#include "shader_recompiler/backend/spirv/emit_spirv.h"
#include "shader_recompiler/compile_stats.h"
#include "shader_recompiler/info.h"
#include "shader_recompiler/ir/abstract_syntax_list.h"
#include "shader_recompiler/ir/basic_block.h"
#include "shader_recompiler/native/capture.h"
#include "shader_recompiler/native/table_plan.h"
#include "shader_recompiler/native/tone_curve.h"
#include "shader_recompiler/ir/ir_emitter.h"
#include "shader_recompiler/ir/passes/ir_passes.h"
#include "shader_recompiler/frontend/control_flow_graph.h"
#include "shader_recompiler/frontend/opcodes.h"
#include "shader_recompiler/ir/passes/srt.h"
#include "core/signals.h"
#include "shader_recompiler/profile.h"
#include "shader_recompiler/recompiler.h"
#include "shader_recompiler/runtime_info.h"
#include "shader_recompiler/frontend/fetch_shader.h"
#include <miniz.h>
#include <set>
#include "shader_recompiler/specialization.h"
#include "common/hash.h"
#include "common/serdes.h"
#include "video_core/renderer_vulkan/vk_graphics_pipeline.h"
#include "video_core/renderer_vulkan/vk_pipeline_serialization.h"

namespace fs = std::filesystem;
using namespace Shader;

extern "C" {
void runtime_blocked_report(void) {}
[[noreturn]] void runtime_thread_exit(void*) { std::abort(); }
int sys_pad_type(void) { return 0; }
int sys_pad_button(uint32_t, int) { return 0; }
}

namespace {

bool g_native = false;
fs::path g_spv_dir;

std::vector<char> ReadAll(const fs::path& path) {
    std::ifstream in(path, std::ios::binary);
    return {std::istreambuf_iterator<char>(in), std::istreambuf_iterator<char>()};
}

struct Census {
    u32 insts = 0, goto_vars = 0, exec_reads = 0, exec_writes = 0, ballots = 0, lane_ops = 0,
        bitcasts = 0, discards = 0;
    u32 ifs = 0, loops = 0, breaks = 0;
};

Census Count(const IR::Program& program) {
    Census c;
    for (const IR::Block* block : program.blocks) {
        for (const IR::Inst& inst : block->Instructions()) {
            ++c.insts;
            switch (inst.GetOpcode()) {
            case IR::Opcode::GetGotoVariable:
            case IR::Opcode::SetGotoVariable:
                ++c.goto_vars;
                break;
            case IR::Opcode::GetExec:
                ++c.exec_reads;
                break;
            case IR::Opcode::SetExec:
                ++c.exec_writes;
                break;
            case IR::Opcode::Ballot:
            case IR::Opcode::InverseBallot:
            case IR::Opcode::BallotFindLsb:
                ++c.ballots;
                break;
            case IR::Opcode::ReadLane:
            case IR::Opcode::ReadFirstLane:
            case IR::Opcode::WriteLane:
            case IR::Opcode::Shuffle:
            case IR::Opcode::ShuffleXor:
            case IR::Opcode::QuadBroadcast:
                ++c.lane_ops;
                break;
            case IR::Opcode::BitCastU16F16:
            case IR::Opcode::BitCastU32F32:
            case IR::Opcode::BitCastF16U16:
            case IR::Opcode::BitCastF32U32:
                ++c.bitcasts;
                break;
            case IR::Opcode::Discard:
            case IR::Opcode::DiscardCond:
                ++c.discards;
                break;
            default:
                break;
            }
        }
    }
    for (const auto& node : program.syntax_list) {
        switch (node.type) {
        case IR::AbstractSyntaxNode::Type::If:
            ++c.ifs;
            break;
        case IR::AbstractSyntaxNode::Type::Loop:
            ++c.loops;
            break;
        case IR::AbstractSyntaxNode::Type::Break:
            ++c.breaks;
            break;
        default:
            break;
        }
    }
    return c;
}

template <class Bytes>
size_t CopyTrailerAt(const Bytes& b) {
    if (b.size() < 8) return b.size();
    u32 count = 0, magic = 0;
    std::memcpy(&count, b.data() + b.size() - 8, 4);
    std::memcpy(&magic, b.data() + b.size() - 4, 4);
    if (magic != Native::CopyTrailerMagic || size_t(count) * 4 + 8 > b.size()) return b.size();
    return b.size() - 8 - size_t(count) * 4;
}

u32* LowWords(size_t count) {
    static u32* block = nullptr;
    static size_t capacity = 0;
    if (count <= capacity) return block;
#ifdef _WIN32
    if (block) VirtualFree(block, 0, MEM_RELEASE);
    block = nullptr;
    capacity = 0;
    const size_t bytes = (count * 4 + 0xFFFF) & ~size_t(0xFFFF);
    for (uintptr_t at = 0x10000000; at < (uintptr_t(1) << 40) && !block; at += 0x10000000) {
        block = static_cast<u32*>(VirtualAlloc(reinterpret_cast<void*>(at), bytes, MEM_RESERVE | MEM_COMMIT, PAGE_READWRITE));
    }
    if (block) capacity = bytes / 4;
#endif
    return block;
}

struct LoadedCapture {
    Native::CaptureHeader header{};
    alignas(RuntimeInfo) unsigned char runtime_storage[sizeof(RuntimeInfo)]{};
    std::array<u32, ShaderParams::NumShaderUserData> user_data{};
    std::vector<u32> code, fetch, flat, copy;
    u32* tess = nullptr;
    std::map<u32, std::vector<u32>> tables;

    RuntimeInfo& Runtime() {
        return *std::launder(reinterpret_cast<RuntimeInfo*>(runtime_storage));
    }
    ShaderParams Params() {
        return ShaderParams{.user_data = user_data, .code = code, .hash = header.hash};
    }

    const char* Load(const fs::path& file) {
        const auto bytes = ReadAll(file);
        auto& h = header;
        if (bytes.size() < sizeof(h)) {
            return "capture cut short";
        }
        std::memcpy(&h, bytes.data(), sizeof(h));
        if (std::memcmp(h.magic, Native::CaptureMagic, 4) != 0 || h.version < 2 ||
            h.version > Native::CaptureVersion || h.runtime_info_size != sizeof(RuntimeInfo)) {
            return "capture of another engine version";
        }
        const char* p = bytes.data() + sizeof(h);
        const char* const end = bytes.data() + bytes.size();
        const auto take = [&](void* out, size_t size) {
            if (end < p + size) {
                return false;
            }
            std::memcpy(out, p, size);
            p += size;
            return true;
        };
        std::vector<u32> ud(h.user_data_count);
        code.resize(h.code_dwords);
        fetch.resize(h.fetch_dwords);
        flat.resize(h.flat_dwords);
        if (!take(runtime_storage, sizeof(RuntimeInfo)) || !take(ud.data(), ud.size() * 4) ||
            !take(code.data(), code.size() * 4) || !take(fetch.data(), fetch.size() * 4) ||
            !take(flat.data(), flat.size() * 4)) {
            return "capture cut short";
        }
        std::memcpy(user_data.data(), ud.data(), std::min(ud.size(), user_data.size()) * 4);
        u32 sharp_count = 0;
        if (h.version >= 3 && !take(&sharp_count, 4)) {
            return "capture cut short";
        }
        std::vector<Native::CaptureSharp> sharps(sharp_count);
        if (!take(sharps.data(), sharps.size() * sizeof(Native::CaptureSharp))) {
            return "capture cut short";
        }
        for (const auto& sharp : sharps) {
            auto& table = tables[sharp.sgpr_base];
            if (table.size() < sharp.dword_offset + 4) {
                table.resize(sharp.dword_offset + 4);
            }
            std::memcpy(&table[sharp.dword_offset], sharp.words, sizeof(sharp.words));
        }
        Native::CaptureTess tess_info{};
        if (h.version >= 4) {
            if (!take(&tess_info, sizeof(tess_info))) {
                return "capture cut short";
            }
            tess = tess_info.size ? LowWords((tess_info.size + 3) / 4) : nullptr;
            if (tess_info.size && !tess) {
                return "no low memory for the tessellation constants";
            }
            if (!take(tess, tess_info.size)) {
                return "capture cut short";
            }
        }
        if (tess_info.size) {
            AmdGpu::Buffer buffer;
            std::memcpy(&buffer, tess_info.words, sizeof(buffer));
            buffer.base_address = reinterpret_cast<uintptr_t>(tess);
            if (buffer.base_address != reinterpret_cast<uintptr_t>(tess)) {
                return "tessellation constants above the V# address range";
            }
            u32 words[4];
            std::memcpy(words, &buffer, sizeof(words));
            if (tess_info.sgpr_base == IR::NumScalarRegs) {
                if (tess_info.dword_offset + 4 <= user_data.size()) {
                    std::memcpy(&user_data[tess_info.dword_offset], words, sizeof(words));
                }
            } else {
                auto& table = tables[tess_info.sgpr_base];
                if (table.size() < tess_info.dword_offset + 4) {
                    table.resize(tess_info.dword_offset + 4);
                }
                std::memcpy(&table[tess_info.dword_offset], words, sizeof(words));
            }
        }
        for (auto& [sgpr, table] : tables) {
            if (sgpr + 1 < user_data.size()) {
                const u32* address = table.data();
                std::memcpy(&user_data[sgpr], &address, sizeof(address));
            }
        }
        if (!fetch.empty() && h.fetch_sgpr + 1 < user_data.size()) {
            const u32* address = fetch.data();
            std::memcpy(&user_data[h.fetch_sgpr], &address, sizeof(address));
        }
        if (const auto at = CopyTrailerAt(bytes); at != bytes.size()) {
            u32 count = 0;
            std::memcpy(&count, bytes.data() + bytes.size() - 8, 4);
            copy.resize(count);
            std::memcpy(copy.data(), bytes.data() + at, size_t(count) * 4);
        }
        if (HwStage(h.hw_stage) == HwStage::Geometry) {
            if (copy.empty()) {
                return "geometry shader without its copy shader (a state of an older engine)";
            }
            Runtime().hw.gs.vs_copy = copy;
        }
        Info::lab_flat = flat.data();
        Info::lab_flat_size = flat.size();
        return nullptr;
    }
};

int WalkNull(const fs::path& file) {
    const auto profile_bytes = ReadAll(file.parent_path() / "profile.bin");
    Profile profile{};
    if (profile_bytes.size() >= 4) {
        std::memcpy(&profile, profile_bytes.data() + 4,
                    std::min<size_t>(profile_bytes.size() - 4, sizeof(Profile)));
    }
    LoadedCapture cap;
    if (const char* error = cap.Load(file)) {
        std::printf("%s: %s\n", file.filename().string().c_str(), error);
        return 2;
    }
    ShaderParams params = cap.Params();
    Info info(HwStage(cap.header.hw_stage), SwStage(cap.header.sw_stage), params);
    Pools pools;
    TranslateProgram(cap.code, pools, info, cap.Runtime(), profile);
    const auto walker = info.srt_info.walker_func;
    const u32 size = info.srt_info.flattened_bufsize_dw;
    if (!walker || size <= NUM_USER_DATA_REGS) {
        std::printf("%s: no walker\n", file.filename().string().c_str());
        return 0;
    }
    std::array<u32, 256> user_data{};
    std::vector<u32> flat(size);
    const auto zeros = [&] {
        for (u32 i = NUM_USER_DATA_REGS; i < size; ++i) {
            if (flat[i] != 0) {
                return false;
            }
        }
        return true;
    };
    std::fill(flat.begin(), flat.end(), 0xCDCDCDCDu);
    walker(user_data.data(), flat.data());
    if (!zeros()) {
        std::printf("%s: FAIL null tables: not all zeros\n", file.filename().string().c_str());
        return 1;
    }
    static std::vector<u32> block(1u << 18);
    const u64 address = reinterpret_cast<u64>(block.data());
    for (u32 i = 0; i + 1 < user_data.size(); i += 2) {
        user_data[i] = u32(address);
        user_data[i + 1] = u32(address >> 32);
    }
    std::fill(flat.begin(), flat.end(), 0xCDCDCDCDu);
    walker(user_data.data(), flat.data());
    if (!zeros()) {
        std::printf("%s: FAIL tables of zeros: not all zeros\n", file.filename().string().c_str());
        return 1;
    }
    std::printf("%s: PASS (%u flattened values)\n", file.filename().string().c_str(),
                size - NUM_USER_DATA_REGS);
    return 0;
}

bool ApplyLabProfile(Profile& profile) {
    const char* set = std::getenv("BB_LAB_PROFILE");
    if (!set) {
        return true;
    }
    const std::map<std::string, bool*> fields = {
        {"needs_manual_interpolation", &profile.needs_manual_interpolation},
        {"uniform_constant_buffers", &profile.uniform_constant_buffers},
        {"uniform_dword_layout", &profile.uniform_dword_layout},
        {"uniform_exact_base", &profile.uniform_exact_base},
        {"flatbuf_storage", &profile.flatbuf_storage},
        {"early_fragment_tests", &profile.early_fragment_tests},
        {"supports_fragment_shader_barycentric", &profile.supports_fragment_shader_barycentric},
        {"supports_amd_shader_explicit_vertex_parameter", &profile.supports_amd_shader_explicit_vertex_parameter},
        {"needs_integer_interpolation_fix", &profile.needs_integer_interpolation_fix},
        {"needs_buffer_offsets", &profile.needs_buffer_offsets},
        {"needs_clip_distance_emulation", &profile.needs_clip_distance_emulation},
        {"supports_shader_stencil_export", &profile.supports_shader_stencil_export},
        {"supports_trinary_minmax", &profile.supports_trinary_minmax},
        {"supports_native_cube_calc", &profile.supports_native_cube_calc},
        {"supports_image_load_store_lod", &profile.supports_image_load_store_lod},
        {"fragment_exact_fp", &profile.fragment_exact_fp},
        {"manual_interpolation_all", &profile.manual_interpolation_all},
    };
    const std::string list = set;
    for (size_t at = 0; at < list.size();) {
        const size_t end = std::min(list.find(',', at), list.size());
        const std::string item = list.substr(at, end - at);
        const size_t eq = item.find('=');
        const auto it = fields.find(item.substr(0, eq));
        if (eq == std::string::npos || it == fields.end()) {
            std::fprintf(stderr, "BB_LAB_PROFILE: unknown %s\n", item.c_str());
            return false;
        }
        *it->second = item[eq + 1] == '1';
        at = end + 1;
    }
    return true;
}

const char* kHeader = "file,stage,hash,gcn_dwords,status,ir_insts,spirv_words,goto_vars,exec_reads,"
                      "exec_writes,ballots,lane_ops,bitcasts,discards,ifs,loops,breaks";

int RunOne(const fs::path& file, const fs::path& profile_file) {
    const auto profile_bytes = ReadAll(profile_file);
    Profile profile{};
    if (profile_bytes.size() < 4) {
        std::printf("%s,,,,no profile.bin\n", file.filename().string().c_str());
        return 2;
    }
    std::memcpy(&profile, profile_bytes.data() + 4,
                std::min<size_t>(profile_bytes.size() - 4, sizeof(Profile)));
    profile.uniform_constant_buffers = g_native;
    profile.uniform_dword_layout = g_native;
    profile.uniform_exact_base = false;
    if (!ApplyLabProfile(profile)) {
        return 2;
    }

    LoadedCapture cap;
    if (const char* error = cap.Load(file)) {
        std::printf("%s,,,,%s\n", file.filename().string().c_str(), error);
        return 2;
    }
    const auto& h = cap.header;
    const auto& code = cap.code;
    RuntimeInfo& runtime_info = cap.Runtime();
    ShaderParams params = cap.Params();
    Info info(HwStage(h.hw_stage), SwStage(h.sw_stage), params);
    Pools pools;
    const auto program = TranslateProgram(code, pools, info, runtime_info, profile);
    Backend::Bindings bindings{};
    const auto spv = Backend::SPIRV::EmitSPIRV(profile, runtime_info, program, bindings);
    if (!g_spv_dir.empty()) {
        std::ofstream out(g_spv_dir / (file.stem().string() + ".spv"), std::ios::binary);
        out.write(reinterpret_cast<const char*>(spv.data()), std::streamsize(spv.size() * 4));
    }
    const Census c = Count(program);
    std::printf("%s,%u,%016llx,%zu,ok,%u,%zu,%u,%u,%u,%u,%u,%u,%u,%u,%u,%u\n",
                file.filename().string().c_str(), h.hw_stage, (unsigned long long)h.hash,
                code.size(), c.insts, spv.size(), c.goto_vars, c.exec_reads, c.exec_writes,
                c.ballots, c.lane_ops, c.bitcasts, c.discards, c.ifs, c.loops, c.breaks);
    std::map<std::string, u32> ops;
    for (const IR::Block* block : program.blocks) {
        for (const IR::Inst& inst : block->Instructions()) {
            ++ops[std::string(IR::NameOf(inst.GetOpcode()))];
        }
    }
    std::printf("OPS|");
    for (const auto& [name, n] : ops) {
        std::printf("%s=%u;", name.c_str(), n);
    }
    std::printf("\n");
    return 0;
}

#ifndef _WIN32
std::string ShellQuote(const std::string& text) {
    std::string quoted = "'";
    for (const char c : text) {
        quoted += c == '\'' ? std::string("'\\''") : std::string(1, c);
    }
    return quoted + "'";
}

std::string RunSelf(const fs::path& self, const std::string& args, int& exit_code) {
    std::string out;
    exit_code = -1;
    const std::string line = ShellQuote(fs::absolute(self).string()) + " " + args;
    if (FILE* pipe = popen(line.c_str(), "r")) {
        char buffer[4096];
        for (size_t got; (got = std::fread(buffer, 1, sizeof buffer, pipe)) > 0;) {
            out.append(buffer, got);
        }
        const int status = pclose(pipe);
        exit_code = status != -1 && WIFEXITED(status) ? WEXITSTATUS(status) : -1;
    }
    return out;
}
#endif

#ifdef _WIN32
std::string RunSelf(const fs::path& self, const std::wstring& args, DWORD& exit_code) {
    static std::mutex spawn_mutex;
    SECURITY_ATTRIBUTES sa{sizeof sa, nullptr, FALSE};
    HANDLE read = nullptr, write = nullptr;
    CreatePipe(&read, &write, &sa, 0);
    STARTUPINFOW si{sizeof si};
    si.dwFlags = STARTF_USESTDHANDLES;
    si.hStdOutput = write;
    si.hStdError = GetStdHandle(STD_ERROR_HANDLE);
    si.hStdInput = GetStdHandle(STD_INPUT_HANDLE);
    std::wstring line = L"\"" + self.wstring() + L"\" " + args;
    PROCESS_INFORMATION pi{};
    std::string out;
    exit_code = ~0u;
    BOOL started = FALSE;
    {
        std::scoped_lock lock{spawn_mutex};
        SetHandleInformation(write, HANDLE_FLAG_INHERIT, HANDLE_FLAG_INHERIT);
        started = CreateProcessW(nullptr, line.data(), nullptr, nullptr, TRUE, CREATE_NO_WINDOW, nullptr,
                                 nullptr, &si, &pi);
        SetHandleInformation(write, HANDLE_FLAG_INHERIT, 0);
    }
    if (started) {
        CloseHandle(write);
        write = nullptr;
        char buffer[4096];
        DWORD got = 0;
        while (ReadFile(read, buffer, sizeof buffer, &got, nullptr) && got) {
            out.append(buffer, got);
        }
        WaitForSingleObject(pi.hProcess, INFINITE);
        GetExitCodeProcess(pi.hProcess, &exit_code);
        CloseHandle(pi.hProcess);
        CloseHandle(pi.hThread);
    }
    if (write) {
        CloseHandle(write);
    }
    CloseHandle(read);
    return out;
}

std::string RunChild(const fs::path& self, const fs::path& file, DWORD& exit_code, std::string& ops) {
    SECURITY_ATTRIBUTES sa{sizeof sa, nullptr, TRUE};
    HANDLE read = nullptr, write = nullptr;
    CreatePipe(&read, &write, &sa, 0);
    SetHandleInformation(read, HANDLE_FLAG_INHERIT, 0);
    STARTUPINFOW si{sizeof si};
    si.dwFlags = STARTF_USESTDHANDLES;
    si.hStdOutput = write;
    si.hStdError = GetStdHandle(STD_ERROR_HANDLE);
    si.hStdInput = GetStdHandle(STD_INPUT_HANDLE);
    std::wstring line = L"\"" + self.wstring() + L"\" --one \"" + file.wstring() + L"\"";
    if (g_native) {
        line += L" --native";
    }
    if (!g_spv_dir.empty()) {
        line += L" --spv-dir \"" + g_spv_dir.wstring() + L"\"";
    }
    PROCESS_INFORMATION pi{};
    std::string out;
    exit_code = ~0u;
    if (CreateProcessW(nullptr, line.data(), nullptr, nullptr, TRUE, CREATE_NO_WINDOW, nullptr,
                       nullptr, &si, &pi)) {
        CloseHandle(write);
        write = nullptr;
        char buffer[4096];
        DWORD got = 0;
        while (ReadFile(read, buffer, sizeof buffer, &got, nullptr) && got) {
            out.append(buffer, got);
        }
        WaitForSingleObject(pi.hProcess, 60000);
        GetExitCodeProcess(pi.hProcess, &exit_code);
        CloseHandle(pi.hProcess);
        CloseHandle(pi.hThread);
    }
    if (write) {
        CloseHandle(write);
    }
    CloseHandle(read);
    std::string csv;
    size_t at = 0;
    while (at < out.size()) {
        const size_t end = out.find('\n', at);
        std::string row = out.substr(at, end == std::string::npos ? std::string::npos : end - at);
        if (!row.empty() && row.back() == '\r') {
            row.pop_back();
        }
        if (row.rfind(file.filename().string(), 0) == 0) {
            csv = row;
        } else if (row.rfind("OPS|", 0) == 0) {
            ops = row.substr(4);
        }
        if (end == std::string::npos) {
            break;
        }
        at = end + 1;
    }
    return csv;
}
#endif

int RunAll(const fs::path& self, const fs::path& dir, const fs::path& report) {
    std::vector<fs::path> files;
    for (const auto& entry : fs::directory_iterator(dir)) {
        if (entry.path().extension() == ".bbs") {
            files.push_back(entry.path());
        }
    }
    std::sort(files.begin(), files.end());
    std::ofstream csv(report);
    csv << kHeader << '\n';
    u32 ok = 0, crashed = 0;
    Census total;
    u32 with_goto = 0, with_exec = 0, with_ballot = 0, with_lane = 0;
    std::map<u32, u32> stages;
    std::map<std::string, std::map<u32, u64>> histogram;
    for (size_t i = 0; i < files.size(); ++i) {
#ifdef _WIN32
        DWORD code = 0;
        std::string ops;
        std::string row = RunChild(self, files[i], code, ops);
#else
        std::string row, ops;
        u32 code = 1;
#endif
        if (row.empty() || code != 0) {
            ++crashed;
            csv << files[i].filename().string() << ",,,,crashed (exit " << code << ")\n";
            continue;
        }
        csv << row << '\n';
        ++ok;
        u32 v[17] = {};
        unsigned long long hash = 0;
        char status[16] = {};
        size_t comma = row.find(',');
        std::sscanf(row.c_str() + comma + 1, "%u,%llx,%u,%15[^,],%u,%u,%u,%u,%u,%u,%u,%u,%u,%u,%u,%u",
                    &v[0], &hash, &v[1], status, &v[2], &v[3], &v[4], &v[5], &v[6], &v[7], &v[8],
                    &v[9], &v[10], &v[11], &v[12], &v[13]);
        ++stages[v[0]];
        for (size_t at = 0; at < ops.size();) {
            const size_t end = ops.find(';', at);
            const std::string pair = ops.substr(at, end == std::string::npos ? std::string::npos : end - at);
            const size_t eq = pair.find('=');
            if (eq != std::string::npos) {
                const u64 n = std::strtoull(pair.c_str() + eq + 1, nullptr, 10);
                histogram[pair.substr(0, eq)][v[0]] += n;
            }
            if (end == std::string::npos) {
                break;
            }
            at = end + 1;
        }
        total.insts += v[2];
        total.goto_vars += v[4];
        total.exec_reads += v[5];
        total.exec_writes += v[6];
        total.ballots += v[7];
        total.lane_ops += v[8];
        total.bitcasts += v[9];
        with_goto += v[4] != 0;
        with_exec += (v[5] + v[6]) != 0;
        with_ballot += v[7] != 0;
        with_lane += v[8] != 0;
        if ((i + 1) % 200 == 0) {
            std::printf("%zu of %zu\n", i + 1, files.size());
        }
    }
    {
        std::ofstream out(fs::path(report).replace_filename(report.stem().string() + "-ops.csv"));
        out << "op,total,fragment,vertex,compute,other\n";
        for (const auto& [name, per_stage] : histogram) {
            unsigned long long total_n = 0, other = 0;
            for (const auto& [stage, n] : per_stage) {
                total_n += n;
                if (stage != 0 && stage != 1 && stage != 6) {
                    other += n;
                }
            }
            const auto at = [&](u32 st) {
                const auto it = per_stage.find(st);
                return it == per_stage.end() ? 0ull : (unsigned long long)it->second;
            };
            out << name << ',' << total_n << ',' << at(0) << ',' << at(1) << ',' << at(6) << ',' << other
                << '\n';
        }
    }
    std::printf("Shader lab: %zu captures, %u compiled, %u crashed the current recompiler\n",
                files.size(), ok, crashed);
    for (const auto& [stage, n] : stages) {
        std::printf("  hardware stage %u: %u\n", stage, n);
    }
    std::printf("  with goto variables: %u (%u in all)\n", with_goto, total.goto_vars);
    std::printf("  with EXEC read or written as data: %u (%u reads, %u writes)\n", with_exec,
                total.exec_reads, total.exec_writes);
    std::printf("  with ballots: %u (%u)\n", with_ballot, total.ballots);
    std::printf("  with lane operations: %u (%u)\n", with_lane, total.lane_ops);
    std::printf("  bit casts: %u in %u IR instructions\n", total.bitcasts, total.insts);
    std::printf("Per shader: %s\n", report.string().c_str());
    return 0;
}

std::vector<std::string> SpvNames(const fs::path& dir);

std::vector<std::string> BatchNames(const fs::path& dir, const fs::path& only) {
    if (only.empty()) {
        return SpvNames(dir);
    }
    std::vector<std::string> names;
    std::ifstream in(only);
    for (std::string line; std::getline(in, line);) {
        if (!line.empty() && line.back() == '\r') line.pop_back();
        if (!line.empty()) names.push_back(line);
    }
    std::sort(names.begin(), names.end());
    return names;
}

std::vector<std::string> SpvNames(const fs::path& dir) {
    std::set<std::string> names;
    std::error_code error;
    for (const auto& entry : fs::directory_iterator(dir, error)) {
        const auto ext = entry.path().extension();
        if (ext == ".spv" || ext == Native::StateExtension) {
            names.insert(entry.path().stem().string());
        }
    }
    return {names.begin(), names.end()};
}

bool WriteAtomically(const fs::path& path, const void* data, size_t size) {
    auto temporary = path;
    temporary += ".tmp";
    {
        std::ofstream out(temporary, std::ios::binary | std::ios::trunc);
        out.write(static_cast<const char*>(data), std::streamsize(size));
        if (!out) {
            return false;
        }
    }
    std::error_code error;
    fs::rename(temporary, path, error);
    return !error;
}

fs::path MetaPath(const fs::path& dir, u64 pgm, size_t perm) {
    char name[40];
    std::snprintf(name, sizeof(name), "0x%016llx.meta", (unsigned long long)HashCombine(pgm, u64(perm)));
    return dir / name;
}

std::vector<u64> KeyStageHashes(const std::vector<char>& key) {
    u32 version = 0, is_compute = 0;
    if (key.size() < 8) return {};
    std::memcpy(&version, key.data(), 4);
    std::memcpy(&is_compute, key.data() + 4, 4);
    const size_t count = is_compute ? 1 : size_t(SwStage::NumLogicalStages);
    if (is_compute > 1 || key.size() < 8 + count * 8) return {};
    std::vector<u64> hashes(count);
    std::memcpy(hashes.data(), key.data() + 8, count * 8);
    return hashes;
}

u32 RemoveOrphanKeys(const fs::path& dir) {
    u32 removed = 0;
    std::error_code error;
    std::vector<fs::path> orphans;
    for (const auto& entry : fs::directory_iterator(dir, error)) {
        if (entry.path().extension() != ".key") continue;
        for (const u64 hash : KeyStageHashes(ReadAll(entry.path()))) {
            char name[40];
            std::snprintf(name, sizeof(name), "0x%016llx.meta", (unsigned long long)hash);
            if (hash && !fs::exists(dir / name, error)) {
                orphans.push_back(entry.path());
                break;
            }
        }
    }
    for (const auto& path : orphans) removed += fs::remove(path, error) ? 1 : 0;
    return removed;
}

void Leave(const fs::path& dir, const std::string& name, u64 pgm, size_t perm) {
    std::error_code error;
    fs::remove(dir / (name + ".spv"), error);
    fs::remove(dir / (name + Native::StateExtension), error);
    fs::remove(MetaPath(dir, pgm, perm), error);
}

int RecompileBatch(const fs::path& dir, size_t start, size_t end, const fs::path& only) {
    const auto profile_bytes = ReadAll(dir / "profile.bin");
    Profile profile{};
    if (profile_bytes.size() != sizeof(Profile)) {
        std::printf("\n@PROFILE\n");
        return 3;
    }
    std::memcpy(&profile, profile_bytes.data(), sizeof(Profile));
    const auto names = BatchNames(dir, only);
    Pools pools;
    end = std::min(end, names.size());
    for (size_t i = start; i < end; ++i) {
        const auto& name = names[i];
        unsigned long long pgm = 0;
        size_t perm = 0;
        const auto out = [&](const char* why) {
            std::printf("\n@OUT %zu %s\n", i, why);
            std::fflush(stdout);
        };
        if (std::sscanf(name.c_str(), "0x%llx_%zu", &pgm, &perm) != 2) {
            out("name");
            continue;
        }
        const auto meta_bytes = ReadAll(MetaPath(dir, pgm, perm));
        StageSpecialization old_spec;
        size_t meta_perm = ~size_t{};
        if (meta_bytes.empty() ||
            !Vulkan::ReadShaderMetaSpec(std::vector<u8>(meta_bytes.begin(), meta_bytes.end()),
                                        old_spec, meta_perm) ||
            meta_perm != perm) {
            out("record of another engine");
            continue;
        }
        LoadedCapture cap;
        if (cap.Load(dir / (name + Native::StateExtension))) {
            out("no state");
            continue;
        }
        std::printf("\n@AT %zu %u %u\n", i, cap.header.hw_stage, cap.header.sw_stage);
        std::fflush(stdout);
        Gcn::fetch_override = cap.fetch.empty() && cap.header.fetch_sgpr != ~0u && old_spec.fetch_shader_data
                                  ? &*old_spec.fetch_shader_data
                                  : nullptr;
        ShaderParams params = cap.Params();
        Info info(HwStage(cap.header.hw_stage), SwStage(cap.header.sw_stage), params);
        RuntimeInfo runtime_info = cap.Runtime();
        const auto program = TranslateProgram(cap.code, pools, info, runtime_info, profile);
        Backend::Bindings bindings = old_spec.start;
        auto spv = Backend::SPIRV::EmitSPIRV(profile, runtime_info, program, bindings);
        const StageSpecialization new_spec(info, runtime_info, profile, old_spec.start);
        old_spec.info = &info;
        if (!(new_spec == old_spec)) {
            std::string why = "specialization changed:";
            if (!(new_spec.runtime_info == old_spec.runtime_info)) why += " runtime";
            if (new_spec.vs_attribs != old_spec.vs_attribs) why += " vs_attribs";
            if (new_spec.fetch_shader_data != old_spec.fetch_shader_data) why += " fetch";
            if (new_spec.fmasks != old_spec.fmasks) why += " fmasks";
            if (new_spec.start != old_spec.start) why += " start";
            if (new_spec.buffers.size() != old_spec.buffers.size() || new_spec.buffers != old_spec.buffers) why += " buffers";
            if (new_spec.images.size() != old_spec.images.size() || new_spec.images != old_spec.images) why += " images";
            if (new_spec.samplers.size() != old_spec.samplers.size() || new_spec.samplers != old_spec.samplers) why += " samplers";
            if (new_spec.bitset != old_spec.bitset) why += " bitset";
            out(why.c_str());
            continue;
        }
        const auto meta = Vulkan::SerializeShaderMeta(info, new_spec, HashCombine(pgm, u64(perm)), perm);
        if (!WriteAtomically(MetaPath(dir, pgm, perm), meta.data(), meta.size()) ||
            !WriteAtomically(dir / (name + ".spv"), spv.data(), spv.size() * sizeof(u32))) {
            out("write failed");
            continue;
        }
        std::printf("\n@OK %zu\n", i);
        std::fflush(stdout);
    }
    std::printf("\n@END\n");
    return 0;
}

std::set<size_t> UsedPerms(const fs::path& cache, u64 pgm);

struct StageRecord {
    u64 pgm = 0;
    size_t perm = 0;
    std::unique_ptr<Info> info;
    StageSpecialization spec{};
    std::optional<Gcn::FetchShaderData> fetch;
};

int FixKeys(const fs::path& dir) {
    const auto profile_bytes = ReadAll(dir / "profile.bin");
    Profile profile{};
    if (profile_bytes.size() != sizeof(Profile)) {
        std::printf("\n@FIXED 0 0 0\n");
        return 3;
    }
    std::memcpy(&profile, profile_bytes.data(), sizeof(Profile));
    std::error_code error;
    std::map<u64, std::unique_ptr<StageRecord>> records;
    const auto record = [&](u64 hash) -> StageRecord* {
        auto& slot = records[hash];
        if (slot) return slot.get();
        char name[40];
        std::snprintf(name, sizeof(name), "0x%016llx.meta", (unsigned long long)hash);
        const auto meta = ReadAll(dir / name);
        if (meta.empty()) return nullptr;
        auto r = std::make_unique<StageRecord>();
        r->info = std::make_unique<Info>();
        r->spec.info = r->info.get();
        Serialization::Archive ar{std::vector<u8>(meta.begin(), meta.end())};
        try {
            if (!Vulkan::LoadShaderMeta(ar, *r->info, r->fetch, r->spec, r->perm, hash)) return nullptr;
        } catch (const std::exception&) {
            return nullptr;
        }
        r->pgm = r->info->pgm_hash;
        slot = std::move(r);
        return slot.get();
    };
    std::map<std::tuple<u64, size_t, u32, u32>, u64> made;
    Pools pools;
    u32 compiled = 0;
    const auto with_start = [&](StageRecord& from, Backend::Bindings start) -> u64 {
        const auto made_key = std::make_tuple(from.pgm, from.perm, start.unified, start.buffer);
        if (const auto it = made.find(made_key); it != made.end()) return it->second;
        start.user_data = from.spec.start.user_data;
        StageSpecialization wanted = from.spec;
        wanted.start = start;
        const auto used = UsedPerms(dir, from.pgm);
        for (const size_t perm : used) {
            StageRecord* other = record(HashCombine(from.pgm, u64(perm)));
            if (other && other->pgm == from.pgm && other->spec == wanted &&
                other->spec.start.unified == start.unified && other->spec.start.buffer == start.buffer) {
                return made[made_key] = HashCombine(from.pgm, u64(perm));
            }
        }
        char state_name[48];
        std::snprintf(state_name, sizeof(state_name), "0x%016llx_%zu", (unsigned long long)from.pgm, from.perm);
        const fs::path state_path = dir / (std::string(state_name) + Native::StateExtension);
        LoadedCapture cap;
        if (cap.Load(state_path)) return made[made_key] = 0;
        Gcn::fetch_override = cap.fetch.empty() && cap.header.fetch_sgpr != ~0u && from.spec.fetch_shader_data
                                  ? &*from.spec.fetch_shader_data
                                  : nullptr;
        ShaderParams params = cap.Params();
        Info info(HwStage(cap.header.hw_stage), SwStage(cap.header.sw_stage), params);
        RuntimeInfo runtime_info = cap.Runtime();
        const auto program = TranslateProgram(cap.code, pools, info, runtime_info, profile);
        Backend::Bindings bindings = start;
        const auto spv = Backend::SPIRV::EmitSPIRV(profile, runtime_info, program, bindings);
        StageSpecialization spec(info, runtime_info, profile, start);
        if (!(spec == wanted)) return made[made_key] = 0;
        size_t perm = used.empty() ? 0 : *used.rbegin() + 1;
        if (perm >= 4096) return made[made_key] = 0;
        const u64 hash = HashCombine(from.pgm, u64(perm));
        char name[48];
        std::snprintf(name, sizeof(name), "0x%016llx_%zu", (unsigned long long)from.pgm, perm);
        const auto meta = Vulkan::SerializeShaderMeta(info, spec, hash, perm);
        const auto state = ReadAll(state_path);
        if (!WriteAtomically(dir / (std::string(name) + Native::StateExtension), state.data(), state.size()) ||
            !WriteAtomically(dir / (std::string(name) + ".spv"), spv.data(), spv.size() * sizeof(u32)) ||
            !WriteAtomically(MetaPath(dir, from.pgm, perm), meta.data(), meta.size())) {
            return made[made_key] = 0;
        }
        ++compiled;
        return made[made_key] = hash;
    };

    std::vector<fs::path> keys;
    for (const auto& entry : fs::directory_iterator(dir, error)) {
        if (entry.path().extension() == ".key") keys.push_back(entry.path());
    }
    std::sort(keys.begin(), keys.end());
    u32 fixed = 0, failed = 0;
    for (const auto& path : keys) {
        const auto bytes = ReadAll(path);
        u32 version = 0, is_compute = 1;
        if (bytes.size() >= 8) {
            std::memcpy(&version, bytes.data(), 4);
            std::memcpy(&is_compute, bytes.data() + 4, 4);
        }
        if (!(version == 5 || version == 6) || is_compute != 0) continue;
        auto hashes = KeyStageHashes(bytes);
        Backend::Bindings expected{};
        bool changed = false, ok = true;
        for (size_t stage = 0; stage < hashes.size() && stage < 5 && ok; ++stage) {
            if (!hashes[stage]) continue;
            StageRecord* r = record(hashes[stage]);
            if (!r) {
                ok = false;
                break;
            }
            for (const auto& image : r->info->images) {
                ok = ok && image.mip_fallback_mode != MipStorageFallbackMode::DynamicIndex;
            }
            if (!ok) break;
            if (r->spec.start.unified != expected.unified || r->spec.start.buffer != expected.buffer) {
                const u64 target = with_start(*r, expected);
                if (!target) {
                    ok = false;
                    break;
                }
                hashes[stage] = target;
                changed = true;
            }
            expected.buffer += u32(r->info->buffers.size());
            expected.unified += u32(r->info->buffers.size() + r->info->images.size() + r->info->samplers.size());
        }
        if (!changed) continue;
        if (!ok) {
            ++failed;
            continue;
        }
        Serialization::Archive ar{std::vector<u8>(bytes.begin() + 8, bytes.end())};
        Vulkan::GraphicsPipelineKey key;
        if (!key.Deserialize(ar)) {
            ++failed;
            continue;
        }
        const std::vector<u8> rest(ar.CurrPtr(), ar.CurrPtr() + ar.Remaining());
        Serialization::Archive check;
        key.Serialize(check);
        if (check.TakeOff().size() + rest.size() + 8 != bytes.size()) {
            ++failed;
            continue;
        }
        for (size_t stage = 0; stage < hashes.size() && stage < key.stage_hashes.size(); ++stage) {
            key.stage_hashes[stage] = hashes[stage];
        }
        Serialization::Archive out;
        key.Serialize(out);
        std::vector<u8> written(bytes.begin(), bytes.begin() + 8);
        const auto key_bytes = out.TakeOff();
        written.insert(written.end(), key_bytes.begin(), key_bytes.end());
        written.insert(written.end(), rest.begin(), rest.end());
        char name[48];
        std::snprintf(name, sizeof(name), "g_0x%016llx.key",
                      (unsigned long long)std::hash<Vulkan::GraphicsPipelineKey>{}(key));
        const fs::path target = dir / name;
        if (!WriteAtomically(target, written.data(), written.size())) {
            ++failed;
            continue;
        }
        if (target != path) fs::remove(path, error);
        ++fixed;
    }
    std::printf("\n@FIXED %u %u %u\n", fixed, compiled, failed);
    return 0;
}

int RecompileCache(const fs::path& self, const fs::path& dir, const fs::path& only = {}) {
    const auto names = BatchNames(dir, only);
    const char* verbose_env = std::getenv("BB_LAB_VERBOSE");
    const bool verbose = verbose_env && verbose_env[0] == '1';
    struct Slice {
        size_t begin = 0, end = 0, kept = 0, left = 0, crashed = 0;
        std::vector<size_t> leave;
        std::map<std::string, u32> reasons;
        bool profile = false;
    };
    const u32 cores = std::thread::hardware_concurrency();
    u32 threads = std::clamp<u32>(cores > 1 ? cores - 1 : 1, 1, 8);
    if (const char* env = std::getenv("BB_RECOMPILE_THREADS"); env && env[0]) {
        threads = u32(std::clamp(std::atoi(env), 1, 32));
    }
    const size_t chunk = std::max<size_t>(32, (names.size() + threads * 4 - 1) / (threads * 4));
    std::vector<Slice> slices;
    for (size_t begin = 0; begin < names.size(); begin += chunk) {
        Slice slice;
        slice.begin = begin;
        slice.end = std::min(names.size(), begin + chunk);
        slices.push_back(std::move(slice));
    }
    const auto run = [&](Slice& slice) {
        size_t next = slice.begin;
        while (next < slice.end) {
#ifdef _WIN32
            DWORD code = 0;
            const std::string output =
                RunSelf(self, L"--recompile-batch \"" + dir.wstring() + L"\" " + std::to_wstring(next) + L" " +
                                  std::to_wstring(slice.end) +
                                  (only.empty() ? std::wstring{} : L" \"" + only.wstring() + L"\""),
                        code);
#else
            int code = 0;
            const std::string output =
                RunSelf(self, "--recompile-batch " + ShellQuote(dir.string()) + " " + std::to_string(next) + " " +
                                  std::to_string(slice.end) +
                                  (only.empty() ? std::string{} : " " + ShellQuote(only.string())),
                        code);
#endif
            if (output.find("@PROFILE") != std::string::npos) {
                slice.profile = true;
                return;
            }
            long long last = -1;
            std::string at_stage;
            bool ended = false;
            size_t at = 0;
            while (at < output.size()) {
                const size_t row_end = output.find('\n', at);
                const std::string row =
                    output.substr(at, row_end == std::string::npos ? std::string::npos : row_end - at);
                size_t index = 0;
                if (std::sscanf(row.c_str(), "@OK %zu", &index) == 1) {
                    ++slice.kept;
                    last = (long long)index;
                } else if (std::sscanf(row.c_str(), "@OUT %zu", &index) == 1) {
                    ++slice.left;
                    slice.leave.push_back(index);
                    last = (long long)index;
                    const auto why = row.find(' ', 5);
                    ++slice.reasons[why == std::string::npos ? "?" : row.substr(why + 1)];
                } else if (row.rfind("@AT ", 0) == 0) {
                    at_stage = row.substr(4);
                } else if (row.rfind("@END", 0) == 0) {
                    ended = true;
                }
                if (row_end == std::string::npos) {
                    break;
                }
                at = row_end + 1;
            }
            if (ended) {
                break;
            }
            const size_t failed = last < 0 ? next : size_t(last) + 1;
            if (failed >= slice.end) {
                break;
            }
            slice.leave.push_back(failed);
            ++slice.crashed;
            ++slice.reasons["compiler refused, hw/sw stage " + at_stage.substr(at_stage.find(' ') + 1)];
            next = failed + 1;
            (void)code;
        }
    };
    std::atomic<size_t> next_slice{0};
    const auto worker = [&] {
        for (size_t i; (i = next_slice.fetch_add(1)) < slices.size();) {
            run(slices[i]);
        }
    };
    std::vector<std::thread> pool;
    for (size_t t = 1; t < std::min<size_t>(threads, slices.size()); ++t) {
        pool.emplace_back(worker);
    }
    worker();
    for (auto& thread : pool) {
        thread.join();
    }
    size_t kept = 0, left = 0, crashed = 0;
    std::vector<size_t> leave;
    std::map<std::string, u32> reasons;
    for (const auto& slice : slices) {
        if (slice.profile) {
            std::printf("Shader cache: written by an engine with another profile; started again\n");
            return 3;
        }
        kept += slice.kept;
        left += slice.left;
        crashed += slice.crashed;
        leave.insert(leave.end(), slice.leave.begin(), slice.leave.end());
        for (const auto& [why, n] : slice.reasons) {
            reasons[why] += n;
        }
    }
    for (const size_t index : leave) {
        unsigned long long pgm = 0;
        size_t perm = 0;
        std::sscanf(names[index].c_str(), "0x%llx_%zu", &pgm, &perm);
        Leave(dir, names[index], pgm, perm);
    }
    const u32 orphan_keys = RemoveOrphanKeys(dir);
    if (verbose) {
        for (const auto& [why, n] : reasons) {
            std::printf("  %u: %s\n", n, why.c_str());
        }
    }
    unsigned fixed = 0, fix_compiled = 0, fix_failed = 0;
    {
#ifdef _WIN32
        DWORD fix_code = 0;
        const std::string fix_out = RunSelf(self, L"--fix-keys \"" + dir.wstring() + L"\"", fix_code);
#else
        int fix_code = 0;
        const std::string fix_out = RunSelf(self, "--fix-keys " + ShellQuote(dir.string()), fix_code);
#endif
        if (const auto at = fix_out.find("@FIXED "); at != std::string::npos) {
            std::sscanf(fix_out.c_str() + at, "@FIXED %u %u %u", &fixed, &fix_compiled, &fix_failed);
        }
        (void)fix_code;
    }
    std::printf("Shader cache: %zu shaders recompiled for this engine, %zu left out (changed or "
                "without a state), %zu the compiler refused, %u pipeline keys without their shaders "
                "removed, %u keys of an older engine made buildable at the start (%u shaders compiled "
                "for them, %u not)\n",
                kept, left, crashed, orphan_keys, fixed, fix_compiled, fix_failed);
    return 0;
}

std::map<u32, char> SpirvBindings(const std::vector<char>& bytes) {
    std::map<u32, char> out;
    if (bytes.size() < 20 || bytes.size() % 4) return out;
    std::vector<u32> w(bytes.size() / 4);
    std::memcpy(w.data(), bytes.data(), bytes.size());
    std::map<u32, u32> binding, set, var_type, ptr_class, ptr_type, elem;
    std::map<u32, char> type_kind;
    for (size_t at = 5; at < w.size();) {
        const u32 count = w[at] >> 16, op = w[at] & 0xffff;
        if (!count || at + count > w.size()) break;
        const u32* a = &w[at + 1];
        if (op == 71 && count >= 4 && a[1] == 33) binding[a[0]] = a[2];
        else if (op == 71 && count >= 4 && a[1] == 34) set[a[0]] = a[2];
        else if (op == 59 && count >= 4) var_type[a[1]] = a[0];
        else if (op == 32 && count >= 4) { ptr_class[a[0]] = a[1]; ptr_type[a[0]] = a[2]; }
        else if (op == 25 && count >= 9) type_kind[a[0]] = a[6] == 2 ? 'W' : 'I';
        else if (op == 26 && count >= 2) type_kind[a[0]] = 'P';
        else if (op == 27 && count >= 3) type_kind[a[0]] = 'I';
        else if ((op == 28 || op == 29) && count >= 3) elem[a[0]] = a[1];
        at += count;
    }
    for (const auto& [var, b] : binding) {
        if (set.count(var) && set[var] != 0) continue;
        const auto vt = var_type.find(var);
        if (vt == var_type.end()) continue;
        const u32 cls = ptr_class[vt->second];
        u32 pointee = ptr_type[vt->second];
        while (elem.count(pointee)) pointee = elem[pointee];
        char kind = '?';
        if (cls == 2) kind = 'U';
        else if (cls == 12) kind = 'S';
        else if (cls == 0 && type_kind.count(pointee)) kind = type_kind[pointee];
        out[b] = kind;
    }
    return out;
}

std::map<u32, std::pair<char, u32>> LayoutBindings(const std::array<const Info*, 5>& stages) {
    std::map<u32, std::pair<char, u32>> out;
    u32 binding = 0;
    for (u32 s = 0; s < stages.size(); ++s) {
        const Info* info = stages[s];
        if (!info) continue;
        for (const auto& buffer : info->buffers) out[binding++] = {buffer.is_uniform ? 'U' : 'S', s};
        for (const auto& image : info->images) {
            const u32 n = image.NumBindings(*info);
            for (u32 i = 0; i < n; ++i) out[binding + i] = {image.is_written ? 'W' : 'I', s};
            binding += n;
        }
        for (size_t i = 0; i < info->samplers.size(); ++i) out[binding++] = {'P', s};
    }
    return out;
}

u32 LayoutMismatches(const std::map<u32, char>& used, const std::map<u32, std::pair<char, u32>>& layout,
                     u32 stage, std::string* first) {
    u32 bad = 0;
    for (const auto& [b, kind] : used) {
        const auto it = layout.find(b);
        const bool ok = it != layout.end() && it->second.first == kind && it->second.second == stage;
        if (!ok) {
            if (first && first->empty()) {
                *first = "stage " + std::to_string(stage) + " binding " + std::to_string(b) + " " + kind +
                         (it == layout.end() ? " not in the layout"
                                             : std::string(" is ") + it->second.first + " of stage " +
                                                   std::to_string(it->second.second));
            }
            ++bad;
        }
    }
    return bad;
}

int CheckCache(const fs::path& dir) {
    std::vector<fs::path> keys;
    std::error_code error;
    for (const auto& entry : fs::directory_iterator(dir, error)) {
        if (entry.path().extension() == ".key") keys.push_back(entry.path());
    }
    std::sort(keys.begin(), keys.end());
    std::map<u64, std::unique_ptr<Info>> infos;
    std::map<std::pair<u64, size_t>, StageSpecialization> loaded;
    std::map<std::string, u32> reasons;
    u32 good = 0;
    std::map<std::pair<u64, size_t>, std::unique_ptr<Info>> own_infos;
    u32 warmup_bad = 0, own_bad = 0, checked = 0;
    std::map<std::string, u32> warmup_examples;
    for (const auto& path : keys) {
        const auto key = ReadAll(path);
        std::array<const Info*, 5> warmup_stages{}, own_stages{};
        std::array<std::map<u32, char>, 5> used{};
        Backend::Bindings expected{};
        bool fs_motion = false, vs_motion = false;
        u32 version = 0, is_compute = 0;
        if (key.size() >= 8) {
            std::memcpy(&version, key.data(), 4);
            std::memcpy(&is_compute, key.data() + 4, 4);
        }
        if (!(version == 5 || version == 6) || is_compute > 1) {
            ++reasons["key version " + std::to_string(version)];
            continue;
        }
        std::string why;
        const auto hashes = KeyStageHashes(key);
        for (size_t stage = 0; stage < hashes.size() && why.empty(); ++stage) {
            const u64 hash = hashes[stage];
            if (!hash) continue;
            char name[40];
            std::snprintf(name, sizeof(name), "0x%016llx.meta", (unsigned long long)hash);
            const auto meta = ReadAll(dir / name);
            if (meta.empty()) {
                why = "no shader record (.meta)";
                break;
            }
            Serialization::Archive ar{std::vector<u8>(meta.begin(), meta.end())};
            auto info = std::make_unique<Info>();
            StageSpecialization spec{};
            spec.info = info.get();
            std::optional<Gcn::FetchShaderData> fetch;
            size_t perm = 0;
            try {
                if (!Vulkan::LoadShaderMeta(ar, *info, fetch, spec, perm, hash)) {
                    Serialization::Archive again{std::vector<u8>(meta.begin(), meta.end())};
                    Serialization::Reader reader{again};
                    u32 meta_version = 0, binary_version = 0;
                    u64 perm_hash = 0;
                    size_t perm_idx = 0;
                    reader.Read(meta_version);
                    reader.Read(binary_version);
                    reader.Read(perm_hash);
                    reader.Read(perm_idx);
                    Info info2;
                    StageSpecialization spec2{};
                    spec2.info = &info2;
                    if (meta_version != 8) {
                        why = "shader record version " + std::to_string(meta_version);
                    } else if (binary_version != 7 &&
                               binary_version != 8) {
                        why = "shader binary version " + std::to_string(binary_version);
                    } else if (!spec2.Deserialize(again)) {
                        why = "specialization unreadable";
                    } else if (!info2.Deserialize(again)) {
                        bool zeros = true;
                        for (size_t i = 0; i < again.Remaining(); ++i) zeros = zeros && again.CurrPtr()[i] == 0;
                        why = std::string("shader info unreadable (srt/table plan), ") +
                              (zeros ? "zero bytes after it" : "other bytes after it");
                    } else if (perm_idx >= 4096 || perm_hash != hash ||
                               HashCombine(info2.pgm_hash, perm_idx) != hash) {
                        why = "hash mismatch";
                    } else {
                        why = "binary version " + std::to_string(binary_version) + " for motion " +
                              std::to_string(info2.hw_stage == HwStage::Vertex &&
                                             spec2.runtime_info.hw.vs.motion_vectors);
                    }
                    break;
                }
            } catch (const std::exception& e) {
                why = std::string("shader record unreadable: ") + e.what();
                break;
            }
            char spv_name[48];
            std::snprintf(spv_name, sizeof(spv_name), "0x%016llx_%zu.spv", (unsigned long long)info->pgm_hash, perm);
            const auto spv = ReadAll(dir / spv_name);
            u32 magic = 0;
            if (spv.size() >= 20) std::memcpy(&magic, spv.data(), 4);
            if (spv.size() < 20 || magic != 0x07230203) {
                why = "no SPIR-V (.spv)";
                break;
            }
            if (u32(info->sw_stage) != (is_compute ? u32(SwStage::Compute) : u32(stage))) {
                why = "shader of another stage";
                break;
            }
            const auto at = std::make_pair(info->pgm_hash, perm);
            if (!is_compute && stage < 5) {
                used[stage] = SpirvBindings(spv);
                auto& own = own_infos[at];
                if (!own) own = std::make_unique<Info>(*info);
                own_stages[stage] = own.get();
                bool mips = false;
                for (const auto& image : info->images)
                    mips |= image.mip_fallback_mode == MipStorageFallbackMode::DynamicIndex;
                if (mips) {
                    why = "an image bound per mip level (built in play)";
                    break;
                }
                if (spec.start.unified != expected.unified || spec.start.buffer != expected.buffer) {
                    why = "a stage compiled after another fragment shader (binding start)";
                    break;
                }
                expected.buffer += u32(info->buffers.size());
                expected.unified += u32(info->buffers.size() + info->images.size() + info->samplers.size());
                if (stage == u32(SwStage::Fragment)) fs_motion = spec.runtime_info.hw.fs.motion_vectors;
                if (stage == u32(SwStage::Vertex))
                    vs_motion = spec.runtime_info.hw_stage == HwStage::Vertex && spec.runtime_info.hw.vs.motion_vectors;
            }
            if (const auto it = loaded.find(at); it != loaded.end()) {
                spec.info = it->second.info;
                if (!(it->second == spec)) {
                    std::string what = "permutation loaded before with another specialization:";
                    if (!(spec.runtime_info == it->second.runtime_info)) what += " runtime";
                    if (spec.vs_attribs != it->second.vs_attribs) what += " vs_attribs";
                    if (spec.fetch_shader_data != it->second.fetch_shader_data) what += " fetch";
                    if (spec.start != it->second.start) what += " start";
                    if (spec.bitset != it->second.bitset) what += " bitset";
                    why = what;
                    break;
                }
            } else {
                auto& owner = infos[info->pgm_hash];
                if (!owner) owner = std::move(info);
                spec.info = owner.get();
                loaded.emplace(at, std::move(spec));
            }
            if (!is_compute && stage < 5) warmup_stages[stage] = infos[at.first].get();
        }
        if (why.empty() && fs_motion && !vs_motion) {
            why = "motion vectors read by the fragment stage, not written by the vertex stage";
        }
        if (why.empty()) {
            ++good;
            if (!is_compute) {
                ++checked;
                const auto warmup_layout = LayoutBindings(warmup_stages);
                const auto own_layout = LayoutBindings(own_stages);
                std::string first;
                u32 w = 0, o = 0;
                for (u32 s = 0; s < 5; ++s) {
                    w += LayoutMismatches(used[s], warmup_layout, s, &first);
                    o += LayoutMismatches(used[s], own_layout, s, nullptr);
                }
                if (w) {
                    ++warmup_bad;
                    if (warmup_examples.size() < 6) ++warmup_examples[path.filename().string() + ": " + first];
                    if (const char* v = std::getenv("BB_LAB_VERBOSE"); v && v[0] == '1' && warmup_bad <= 4) {
                        std::string text = path.filename().string() + "\n    layout:";
                        for (const auto& [b, ks] : warmup_layout)
                            text += " " + std::to_string(b) + ks.first + std::to_string(ks.second);
                        for (u32 s = 0; s < 5; ++s) {
                            if (!warmup_stages[s]) continue;
                            const Info& i = *warmup_stages[s];
                            char head[160];
                            std::snprintf(head, sizeof(head), "\n    stage %u pgm %#llx: %zu buffers, %zu images, %zu samplers; uses:",
                                          s, (unsigned long long)i.pgm_hash, i.buffers.size(), i.images.size(), i.samplers.size());
                            text += head;
                            for (const auto& [b, k] : used[s]) text += " " + std::to_string(b) + k;
                        }
                        std::printf("%s\n", text.c_str());
                    }
                }
                if (o) ++own_bad;
            }
        } else {
            ++reasons[why];
        }
    }
    std::printf("Cache check: %u of %zu pipeline keys would be built at the start\n", good, keys.size());
    for (const auto& [why, n] : reasons) {
        std::printf("  %u: %s\n", n, why.c_str());
    }
    std::printf("Layout check: %u of %u graphics pipelines have a layout that does not match their "
                "shaders as the warm-up builds them; %u with each permutation's own shader info\n",
                warmup_bad, checked, own_bad);
    for (const auto& [example, n] : warmup_examples) {
        std::printf("  %s\n", example.c_str());
    }
    return warmup_bad ? 4 : 0;
}

int Bench(const fs::path& dir, int rounds, const fs::path& report) {
    const auto profile_bytes = ReadAll(dir / "profile.bin");
    Profile profile{};
    if (profile_bytes.size() != sizeof(Profile)) {
        std::fprintf(stderr, "%s: profile of another engine\n", (dir / "profile.bin").string().c_str());
        return 3;
    }
    std::memcpy(&profile, profile_bytes.data(), sizeof(Profile));
    if (!ApplyLabProfile(profile)) {
        return 2;
    }
    struct Entry {
        std::string name;
        LoadedCapture cap;
        StageSpecialization spec;
        u32 stage = 0;
        size_t gcn_dwords = 0;
        std::array<unsigned long long, CompileStats::Count> best{};
        unsigned long long best_total = ~0ull;
        size_t ir_insts = 0, spv_words = 0;
        u64 spv_hash = 0;
        bool spec_same = false;
    };
    std::vector<std::unique_ptr<Entry>> entries;
    for (const auto& name : SpvNames(dir)) {
        unsigned long long pgm = 0;
        size_t perm = 0;
        if (std::sscanf(name.c_str(), "0x%llx_%zu", &pgm, &perm) != 2) {
            continue;
        }
        auto e = std::make_unique<Entry>();
        size_t meta_perm = ~size_t{};
        const auto meta = ReadAll(MetaPath(dir, pgm, perm));
        if (meta.empty() ||
            !Vulkan::ReadShaderMetaSpec(std::vector<u8>(meta.begin(), meta.end()), e->spec, meta_perm) ||
            meta_perm != perm || e->cap.Load(dir / (name + Native::StateExtension)) ||
            e->cap.code.empty()) {
            continue;
        }
        e->name = name;
        e->stage = e->cap.header.hw_stage;
        e->gcn_dwords = e->cap.code.size();
        e->cap = {};
        entries.push_back(std::move(e));
    }
    CompileStats::last_report_ns = 1ll << 62;
    Pools pools;
    for (int round = 0; round < rounds; ++round) {
        for (auto& e : entries) {
            auto& cap = e->cap;
            cap = {};
            cap.Load(dir / (e->name + Native::StateExtension));
            Gcn::fetch_override = cap.fetch.empty() && cap.header.fetch_sgpr != ~0u && e->spec.fetch_shader_data
                                      ? &*e->spec.fetch_shader_data
                                      : nullptr;
            std::array<unsigned long long, CompileStats::Count> before{};
            for (int p = 0; p < CompileStats::Count; ++p) {
                before[p] = CompileStats::phase_ns[p].load();
            }
            ShaderParams params = cap.Params();
            Info info(HwStage(cap.header.hw_stage), SwStage(cap.header.sw_stage), params);
            RuntimeInfo runtime_info = cap.Runtime();
            const auto program = TranslateProgram(cap.code, pools, info, runtime_info, profile);
            Backend::Bindings bindings = e->spec.start;
            const auto spv = Backend::SPIRV::EmitSPIRV(profile, runtime_info, program, bindings);
            std::array<unsigned long long, CompileStats::Count> took{};
            unsigned long long total = 0;
            for (int p = 0; p < CompileStats::Count; ++p) {
                took[p] = CompileStats::phase_ns[p].load() - before[p];
                total += took[p];
            }
            if (total < e->best_total) {
                e->best_total = total;
                e->best = took;
            }
            if (round == 0) {
                for (const IR::Block* block : program.blocks) {
                    e->ir_insts += block->Instructions().size();
                }
                e->spv_words = spv.size();
                e->spv_hash = 0xcbf29ce484222325ull;
                for (const u32 w : spv) {
                    e->spv_hash = (e->spv_hash ^ w) * 0x100000001b3ull;
                }
                const StageSpecialization new_spec(info, runtime_info, profile, e->spec.start);
                e->spec.info = &info;
                e->spec_same = new_spec == e->spec;
                if (!g_spv_dir.empty()) {
                    std::ofstream out(g_spv_dir / (e->name + ".spv"), std::ios::binary);
                    out.write(reinterpret_cast<const char*>(spv.data()), std::streamsize(spv.size() * 4));
                }
            }
        }
    }
    Gcn::fetch_override = nullptr;
    static constexpr const char* kPhase[] = {"decode", "translate", "ssa", "resources", "passes",
                                             "structurize", "late", "spirv"};
    std::ofstream csv(report);
    csv << "name,stage,gcn_dwords,ir_insts,spv_words,spv_hash,spec_same,total_us";
    for (const char* p : kPhase) csv << ',' << p << "_us";
    csv << '\n';
    std::array<double, CompileStats::Count> sum{};
    double all = 0;
    size_t words = 0, insts = 0, changed_spec = 0;
    std::vector<std::pair<double, std::string>> slowest;
    for (const auto& e : entries) {
        char line[256];
        std::snprintf(line, sizeof(line), "%s,%u,%zu,%zu,%zu,%016llx,%d,%.1f", e->name.c_str(),
                      e->stage, e->gcn_dwords, e->ir_insts, e->spv_words,
                      (unsigned long long)e->spv_hash, e->spec_same ? 1 : 0, e->best_total / 1e3);
        csv << line;
        for (int p = 0; p < CompileStats::Count; ++p) {
            csv << ',' << e->best[p] / 1e3;
            sum[p] += e->best[p] / 1e6;
        }
        csv << '\n';
        all += e->best_total / 1e6;
        words += e->spv_words;
        insts += e->ir_insts;
        changed_spec += e->spec_same ? 0 : 1;
        slowest.push_back({e->best_total / 1e6, e->name});
    }
    std::sort(slowest.rbegin(), slowest.rend());
    std::printf("Bench: %zu shaders, best of %d rounds: %.1f ms (%.2f avg)\n", entries.size(), rounds,
                all, entries.empty() ? 0.0 : all / entries.size());
    std::printf(" ");
    for (int p = 0; p < CompileStats::Count; ++p) {
        std::printf(" %s %.1f", kPhase[p], sum[p]);
    }
    std::printf(" ms\n  IR instructions %zu, SPIR-V words %zu, specializations changed %zu\n", insts,
                words, changed_spec);
    std::printf("  slowest:");
    for (size_t i = 0; i < std::min<size_t>(5, slowest.size()); ++i) {
        std::printf(" %s %.1f ms;", slowest[i].second.c_str(), slowest[i].first);
    }
    std::printf("\nPer shader: %s\n", report.string().c_str());
    return 0;
}

std::vector<u8> ReadBytes(const fs::path& path) {
    const auto v = ReadAll(path);
    return {v.begin(), v.end()};
}

bool WriteBytes(const fs::path& path, const std::vector<u8>& bytes) {
    return WriteAtomically(path, bytes.data(), bytes.size());
}

struct StateLayout {
    Native::CaptureHeader header{};
    size_t code_at = 0, fetch_at = 0, rest_at = 0;
    bool Parse(const std::vector<u8>& b) {
        if (b.size() < sizeof(header)) return false;
        std::memcpy(&header, b.data(), sizeof(header));
        if (std::memcmp(header.magic, Native::CaptureMagic, 4) != 0) return false;
        code_at = sizeof(header) + header.runtime_info_size + size_t(header.user_data_count) * 4;
        fetch_at = code_at + size_t(header.code_dwords) * 4;
        rest_at = fetch_at + size_t(header.fetch_dwords) * 4;
        return rest_at <= b.size();
    }
};

int ExportCatalog(const fs::path& cache, const fs::path& out) {
    std::error_code error;
    fs::create_directories(out, error);
    u32 states = 0, metas = 0, keys = 0, dead_keys = 0;
    std::string listed;
    const auto has_meta = [&](u64 hash) {
        char name[40];
        std::snprintf(name, sizeof(name), "0x%016llx.meta", (unsigned long long)hash);
        return fs::exists(cache / name, error);
    };
    for (const auto& entry : fs::directory_iterator(cache, error)) {
        const auto& path = entry.path();
        const auto ext = path.extension();
        if (ext == ".key") {
            const auto hashes = KeyStageHashes(ReadAll(path));
            if (hashes.empty() || std::ranges::any_of(hashes, [&](u64 h) { return h && !has_meta(h); })) {
                ++dead_keys;
                continue;
            }
        }
        if (ext == Native::StateExtension) {
            unsigned long long pgm = 0;
            size_t perm = 0;
            if (std::sscanf(path.stem().string().c_str(), "0x%llx_%zu", &pgm, &perm) != 2 ||
                !fs::exists(MetaPath(cache, pgm, perm), error)) {
                continue;
            }
        }
        if (ext == ".key" || ext == ".meta" || path.filename() == "profile.bin") {
            fs::copy_file(path, out / path.filename(), fs::copy_options::overwrite_existing, error);
            listed += path.filename().string() + "\n";
            ext == ".key" ? ++keys : ext == ".meta" ? ++metas : 0;
        } else if (ext == Native::StateExtension) {
            auto bytes = ReadBytes(path);
            StateLayout layout;
            if (!layout.Parse(bytes)) continue;
            std::vector<u8> stripped(bytes.begin(), bytes.begin() + layout.code_at);
            stripped.insert(stripped.end(), bytes.begin() + layout.rest_at, bytes.begin() + CopyTrailerAt(bytes));
            auto header = layout.header;
            header.code_dwords = 0;
            header.fetch_dwords = 0;
            std::memcpy(stripped.data(), &header, sizeof(header));
            if (WriteBytes(out / path.filename(), stripped)) {
                ++states;
                listed += path.filename().string() + "\n";
            }
        }
    }
    auto driver_cache = cache;
    driver_cache += ".vkcache";
    if (const auto blob = ReadBytes(driver_cache); blob.size() >= 16 && std::memcmp(blob.data(), "BPBC", 4) == 0) {
        u32 vendor = 0, device = 0;
        std::memcpy(&vendor, blob.data() + 8, 4);
        std::memcpy(&device, blob.data() + 12, 4);
        char text[64];
        const int size = std::snprintf(text, sizeof(text), "vendor %04X device %04X\n", vendor, device);
        WriteAtomically(out / "gpu.txt", text, size_t(size));
    }
    WriteAtomically(out / "files.txt", listed.data(), listed.size());
    std::printf("Shader catalog: %u states, %u specializations, %u pipeline keys in %s (%u keys without "
                "their shaders left out)\n",
                states, metas, keys, out.string().c_str(), dead_keys);
    return 0;
}

struct CodeIndex {
    std::vector<std::vector<u8>> blobs;
    std::map<u64, std::pair<size_t, size_t>> programs;
    std::map<u64, std::pair<size_t, size_t>> copies;

    static std::vector<u8> Undcx(std::vector<u8>&& data) {
        if (data.size() < 0x4c || std::memcmp(data.data(), "DCX", 4) != 0) return std::move(data);
        if (std::memcmp(data.data() + 0x28, "DFLT", 4) != 0) return {};
        const auto be32 = [&](size_t at) {
            return u32(data[at]) << 24 | u32(data[at + 1]) << 16 | u32(data[at + 2]) << 8 | data[at + 3];
        };
        const auto dca = std::search(data.begin(), data.end(), "DCA", "DCA" + 4) - data.begin();
        if (size_t(dca) + 8 > data.size()) return {};
        const size_t skip = be32(size_t(dca) + 4);
        mz_ulong size = be32(0x1c);
        std::vector<u8> out(size);
        if (mz_uncompress(out.data(), &size, data.data() + dca + skip, mz_ulong(data.size() - dca - skip)) != MZ_OK) {
            return {};
        }
        out.resize(size);
        return out;
    }

    void Add(std::vector<u8>&& blob) {
        const size_t index = blobs.size();
        static constexpr char kSig[] = "OrbShdr";
        for (size_t at = 0; (at = std::string_view(reinterpret_cast<const char*>(blob.data()), blob.size()).find(kSig, at)) != std::string_view::npos; ++at) {
            if (at + 28 > blob.size()) break;
            u64 hash;
            std::memcpy(&hash, blob.data() + at + 16, 8);
            for (size_t p = at & ~size_t(3); p >= 8 && at - p < 0x40000; p -= 4) {
                u32 token, imm;
                std::memcpy(&token, blob.data() + p, 4);
                std::memcpy(&imm, blob.data() + p + 4, 4);
                if (token == 0xBEEB03FFu && p + (size_t(imm) + 1) * 8 == at) {
                    u32 bits;
                    std::memcpy(&bits, blob.data() + at + 8, 4);
                    ((bits >> 2 & 15) == 6 ? copies : programs).try_emplace(hash, index, p);
                    break;
                }
            }
        }
        blobs.push_back(std::move(blob));
    }

    std::vector<u32> Code(u64 hash, bool copy_shader = false) const {
        auto it = programs.find(hash);
        if (copy_shader) {
            if (const auto copy = copies.find(hash); copy != copies.end()) it = copy;
        }
        if (it == programs.end()) return {};
        const auto& blob = blobs[it->second.first];
        const size_t start = it->second.second;
        u32 imm;
        std::memcpy(&imm, blob.data() + start + 4, 4);
        const size_t footer = start + (size_t(imm) + 1) * 8;
        u32 bits;
        std::memcpy(&bits, blob.data() + footer + 8, 4);
        const size_t length = bits >> 8;
        if (start + length > blob.size()) return {};
        std::vector<u32> code(length / 4);
        std::memcpy(code.data(), blob.data() + start, code.size() * 4);
        return code;
    }
};

std::set<size_t> UsedPerms(const fs::path& cache, u64 pgm) {
    std::set<size_t> used;
    char prefix[32];
    std::snprintf(prefix, sizeof(prefix), "0x%016llx_", (unsigned long long)pgm);
    std::error_code error;
    for (const auto& entry : fs::directory_iterator(cache, error)) {
        const auto name = entry.path().stem().string();
        if (name.rfind(prefix, 0) == 0) used.insert(std::strtoull(name.c_str() + std::strlen(prefix), nullptr, 10));
    }
    return used;
}

int SeedCache(const fs::path& self, const fs::path& catalog, const fs::path& cache, const fs::path& game,
              const fs::path& eboot) {
    std::error_code error;
    const auto catalog_profile = ReadBytes(catalog / "profile.bin");
    const auto cache_profile = ReadBytes(cache / "profile.bin");
    static constexpr size_t GpuBytes = offsetof(Profile, uniform_constant_buffers);
    if (catalog_profile.size() < GpuBytes || cache_profile.size() < GpuBytes) {
        std::printf("Shader catalog: made on another kind of GPU than this one; not used\n");
        return 5;
    }
    if (std::memcmp(catalog_profile.data(), cache_profile.data(), GpuBytes) != 0) {
        unsigned catalog_vendor = 0;
        if (const auto text = ReadBytes(catalog / "gpu.txt"); !text.empty()) {
            std::sscanf(std::string(text.begin(), text.end()).c_str(), "vendor %x", &catalog_vendor);
        }
        u32 cache_vendor = 0;
        auto driver_cache = cache;
        driver_cache += ".vkcache";
        if (const auto blob = ReadBytes(driver_cache); blob.size() >= 16 && std::memcmp(blob.data(), "BPBC", 4) == 0) {
            std::memcpy(&cache_vendor, blob.data() + 8, 4);
        }
        if (catalog_vendor == 0 || catalog_vendor != cache_vendor) {
            std::printf("Shader catalog: made on another kind of GPU than this one; not used\n");
            return 5;
        }
        std::printf("Shader catalog: made on another GPU of this vendor (%04X); each shader compiled for this one\n",
                    catalog_vendor);
    }
    const bool other_switches = catalog_profile != cache_profile;
    std::set<std::string> own;
    if (std::ifstream list{catalog / "files.txt"}) {
        for (std::string line; std::getline(list, line);) {
            if (!line.empty() && line.back() == '\r') line.pop_back();
            if (!line.empty()) own.insert(line);
        }
    }
    const auto in_catalog = [&](const fs::path& path) { return own.empty() || own.count(path.filename().string()) != 0; };
    CodeIndex index;
    for (const auto& entry : fs::directory_iterator(game / "dvdroot_ps4" / "shader", error)) {
        if (entry.path().extension() == ".dcx") index.Add(CodeIndex::Undcx(ReadBytes(entry.path())));
    }
    index.Add(ReadBytes(eboot));
    std::map<u64, u64> renamed_hashes;
    std::vector<std::string> added_names;
    u32 added = 0, missing = 0, present = 0, renumbered = 0, unreadable = 0;
    for (const auto& entry : fs::directory_iterator(catalog, error)) {
        if (entry.path().extension() != Native::StateExtension || !in_catalog(entry.path())) continue;
        const auto name = entry.path().stem().string();
        unsigned long long pgm = 0;
        size_t perm = 0;
        if (std::sscanf(name.c_str(), "0x%llx_%zu", &pgm, &perm) != 2) continue;
        auto bytes = ReadBytes(entry.path());
        StateLayout layout;
        auto meta = ReadBytes(MetaPath(catalog, pgm, perm));
        if (!layout.Parse(bytes) || meta.size() < 24) {
            ++unreadable;
            continue;
        }
        if (u32 version; std::memcpy(&version, meta.data(), 4), version == 7) {
            version = 8;
            std::memcpy(meta.data(), &version, 4);
        }
        const auto code = index.Code(pgm);
        if (code.empty()) {
            ++missing;
            continue;
        }
        const auto used = UsedPerms(cache, pgm);
        StageSpecialization catalog_spec;
        size_t catalog_perm = 0;
        if (!Vulkan::ReadShaderMetaSpec(std::vector<u8>(meta), catalog_spec, catalog_perm)) {
            ++unreadable;
            continue;
        }
        const Info valid_marker{};
        catalog_spec.info = &valid_marker;
        bool have = false;
        for (const size_t mine : used) {
            StageSpecialization spec;
            size_t mine_perm = 0;
            const auto mine_meta = ReadBytes(MetaPath(cache, pgm, mine));
            if (!mine_meta.empty() && Vulkan::ReadShaderMetaSpec(std::vector<u8>(mine_meta), spec, mine_perm)) {
                spec.info = &valid_marker;
                if (spec == catalog_spec) {
                    have = true;
                    renamed_hashes[HashCombine(u64(pgm), u64(perm))] = HashCombine(u64(pgm), u64(mine));
                    break;
                }
            }
        }
        if (have) {
            ++present;
            continue;
        }
        size_t target = perm;
        if (used.count(perm)) {
            ++renumbered;
            target = *used.rbegin() + 1;
            while (used.count(target)) ++target;
        }
        if (target >= 4096) continue;
        const u64 old_hash = HashCombine(u64(pgm), u64(perm));
        const u64 new_hash = HashCombine(u64(pgm), u64(target));
        renamed_hashes[old_hash] = new_hash;
        auto header = layout.header;
        header.code_dwords = u32(code.size());
        std::vector<u8> full(bytes.begin(), bytes.begin() + layout.code_at);
        full.insert(full.end(), reinterpret_cast<const u8*>(code.data()),
                    reinterpret_cast<const u8*>(code.data() + code.size()));
        full.insert(full.end(), bytes.begin() + layout.code_at, bytes.end());
        std::memcpy(full.data(), &header, sizeof(header));
        if (HwStage(header.hw_stage) == HwStage::Geometry && CopyTrailerAt(bytes) == bytes.size()) {
            alignas(RuntimeInfo) unsigned char runtime_storage[sizeof(RuntimeInfo)];
            std::memcpy(runtime_storage, bytes.data() + sizeof(header), sizeof(RuntimeInfo));
            const auto& runtime = *std::launder(reinterpret_cast<const RuntimeInfo*>(runtime_storage));
            const auto copy = index.Code(runtime.hw.gs.vs_copy_hash, true);
            if (copy.empty()) {
                ++missing;
                continue;
            }
            const u32 trailer[2] = {u32(copy.size()), Native::CopyTrailerMagic};
            full.insert(full.end(), reinterpret_cast<const u8*>(copy.data()),
                        reinterpret_cast<const u8*>(copy.data() + copy.size()));
            full.insert(full.end(), reinterpret_cast<const u8*>(trailer), reinterpret_cast<const u8*>(trailer + 2));
        }
        char target_name[48];
        std::snprintf(target_name, sizeof(target_name), "0x%016llx_%zu", pgm, target);
        std::memcpy(meta.data() + 8, &new_hash, 8);
        const u64 target64 = target;
        std::memcpy(meta.data() + 16, &target64, 8);
        if (WriteBytes(cache / (std::string(target_name) + Native::StateExtension), full) &&
            WriteBytes(MetaPath(cache, pgm, target), meta)) {
            ++added;
            added_names.push_back(target_name);
        }
    }
    u32 keys = 0;
    for (const auto& entry : fs::directory_iterator(catalog, error)) {
        if (entry.path().extension() != ".key" || !in_catalog(entry.path()) ||
            fs::exists(cache / entry.path().filename(), error)) {
            continue;
        }
        auto key = ReadBytes(entry.path());
        for (const auto& [from, to] : renamed_hashes) {
            if (from == to) continue;
            for (size_t at = 0; at + 8 <= key.size(); at += 1) {
                u64 v;
                std::memcpy(&v, key.data() + at, 8);
                if (v == from) std::memcpy(key.data() + at, &to, 8);
            }
        }
        if (WriteBytes(cache / entry.path().filename(), key)) ++keys;
    }
    std::printf("Shader catalog: %u permutations added (%u renumbered), %u here already, %u not in this "
                "game's files, %u records of another engine, %u pipeline keys; compiling them%s\n",
                added, renumbered, present, missing, unreadable, keys,
                other_switches ? " (made with other compile switches: each one checked)" : "");
    if (added_names.empty()) {
        return 0;
    }
    const fs::path list = cache / "catalog-seed.txt";
    {
        std::ofstream out(list, std::ios::trunc);
        for (const auto& name : added_names) out << name << '\n';
    }
    const int code = RecompileCache(self, cache, list);
    fs::remove(list, error);
    return code;
}

Profile CaptureProfile(const fs::path& file) {
    const auto bytes = ReadAll(file.parent_path() / "profile.bin");
    Profile profile{};
    if (bytes.size() >= 4) {
        std::memcpy(&profile, bytes.data() + 4, std::min<size_t>(bytes.size() - 4, sizeof(Profile)));
    }
    profile.uniform_constant_buffers = g_native;
    profile.uniform_dword_layout = g_native;
    profile.uniform_exact_base = false;
    return profile;
}

std::vector<int> Dominators(u32 count, u32 entry, const std::vector<std::vector<u32>>& out,
                            const std::vector<std::vector<u32>>& in) {
    std::vector<u32> post;
    std::vector<bool> seen(count);
    std::vector<std::pair<u32, size_t>> stack{{entry, 0}};
    seen[entry] = true;
    while (!stack.empty()) {
        const u32 v = stack.back().first;
        if (stack.back().second < out[v].size()) {
            const u32 w = out[v][stack.back().second++];
            if (!seen[w]) {
                seen[w] = true;
                stack.push_back({w, 0});
            }
        } else {
            post.push_back(v);
            stack.pop_back();
        }
    }
    std::vector<int> order(count, -1);
    for (size_t i = 0; i < post.size(); ++i) {
        order[post[i]] = int(post.size() - 1 - i);
    }
    std::vector<int> idom(count, -1);
    idom[entry] = int(entry);
    for (bool changed = true; changed;) {
        changed = false;
        for (auto it = post.rbegin(); it != post.rend(); ++it) {
            const u32 v = *it;
            if (v == entry) {
                continue;
            }
            int dom = -1;
            for (const u32 p : in[v]) {
                if (idom[p] < 0) {
                    continue;
                }
                if (dom < 0) {
                    dom = int(p);
                    continue;
                }
                int a = int(p), b = dom;
                while (a != b) {
                    while (order[a] > order[b]) a = idom[a];
                    while (order[b] > order[a]) b = idom[b];
                }
                dom = a;
            }
            if (dom >= 0 && idom[v] != dom) {
                idom[v] = dom;
                changed = true;
            }
        }
    }
    return idom;
}

struct FlowShape {
    u32 blocks = 0, branches = 0, exec_branches = 0, exits = 0, loops = 0, irreducible = 0,
        multi_exit_loops = 0, side_entries = 0, max_loop_depth = 0;
};

FlowShape AnalyzeFlow(std::span<const Gcn::GcnInst> insts) {
    FlowShape shape;
    Common::ObjectPool<Gcn::Block> pool{64};
    Gcn::CFG cfg{pool, insts};
    std::vector<const Gcn::Block*> blocks;
    std::map<const Gcn::Block*, u32> index;
    for (const auto& block : cfg) {
        index[&block] = u32(blocks.size());
        blocks.push_back(&block);
    }
    const u32 n = u32(blocks.size());
    shape.blocks = n;
    if (n == 0) {
        return shape;
    }
    std::vector<std::vector<u32>> succ(n + 1), pred(n + 1);
    const auto link = [&](u32 a, u32 b) {
        succ[a].push_back(b);
        pred[b].push_back(a);
    };
    for (u32 i = 0; i < n; ++i) {
        const auto* block = blocks[i];
        if (block->end_class == Gcn::EndClass::Exit) {
            ++shape.exits;
            link(i, n);
            continue;
        }
        if (block->branch_true) link(i, index.at(block->branch_true));
        if (block->branch_false && block->branch_false != block->branch_true) {
            link(i, index.at(block->branch_false));
        }
        if (succ[i].size() == 2) {
            ++shape.branches;
            const auto op = block->end_inst.opcode;
            shape.exec_branches += op == Gcn::Opcode::S_CBRANCH_EXECZ || op == Gcn::Opcode::S_CBRANCH_EXECNZ;
        }
    }
    const auto idom = Dominators(n + 1, 0, succ, pred);
    const auto dominates = [&](u32 a, u32 b) {
        for (int v = int(b); v >= 0;) {
            if (u32(v) == a) return true;
            if (idom[v] == v) return false;
            v = idom[v];
        }
        return false;
    };
    std::set<std::pair<u32, u32>> back;
    std::map<u32, std::vector<u32>> latches;
    {
        std::vector<u8> state(n + 1);
        std::vector<std::pair<u32, size_t>> stack{{0, 0}};
        state[0] = 1;
        while (!stack.empty()) {
            const u32 v = stack.back().first;
            if (stack.back().second < succ[v].size()) {
                const u32 w = succ[v][stack.back().second++];
                if (state[w] == 0) {
                    state[w] = 1;
                    stack.push_back({w, 0});
                } else if (state[w] == 1) {
                    if (dominates(w, v)) {
                        back.insert({v, w});
                        latches[w].push_back(v);
                    } else {
                        ++shape.irreducible;
                    }
                }
            } else {
                state[v] = 2;
                stack.pop_back();
            }
        }
    }
    std::vector<std::vector<bool>> bodies;
    for (const auto& [header, from] : latches) {
        std::vector<bool> body(n + 1);
        body[header] = true;
        std::vector<u32> work(from.begin(), from.end());
        while (!work.empty()) {
            const u32 v = work.back();
            work.pop_back();
            if (body[v]) continue;
            body[v] = true;
            for (const u32 p : pred[v]) work.push_back(p);
        }
        std::set<u32> exits;
        for (u32 v = 0; v < n; ++v) {
            if (!body[v]) continue;
            for (const u32 w : succ[v]) {
                if (!body[w]) exits.insert(w);
            }
        }
        ++shape.loops;
        shape.multi_exit_loops += exits.size() > 1;
        bodies.push_back(std::move(body));
    }
    std::vector<int> innermost(n + 1, -1);
    std::vector<u32> loop_depth(n + 1, 0);
    for (size_t l = 0; l < bodies.size(); ++l) {
        const auto size = std::count(bodies[l].begin(), bodies[l].end(), true);
        for (u32 v = 0; v < n; ++v) {
            if (!bodies[l][v]) continue;
            shape.max_loop_depth = std::max(shape.max_loop_depth, ++loop_depth[v]);
            if (innermost[v] < 0 ||
                size < std::count(bodies[innermost[v]].begin(), bodies[innermost[v]].end(), true)) {
                innermost[v] = int(l);
            }
        }
    }
    const auto ipdom = Dominators(n + 1, n, pred, succ);
    for (u32 b = 0; b < n; ++b) {
        if (succ[b].size() != 2 || back.contains({b, succ[b][0]}) || back.contains({b, succ[b][1]})) {
            continue;
        }
        const int join = ipdom[b] < 0 ? int(n) : ipdom[b];
        const auto* loop = innermost[b] >= 0 ? &bodies[innermost[b]] : nullptr;
        std::vector<bool> region(n + 1);
        std::vector<u32> work(succ[b].begin(), succ[b].end());
        while (!work.empty()) {
            const u32 v = work.back();
            work.pop_back();
            if (v == n || int(v) == join || region[v] || v == b || (loop && !(*loop)[v])) continue;
            region[v] = true;
            for (const u32 w : succ[v]) {
                if (!back.contains({v, w})) work.push_back(w);
            }
        }
        bool side_entry = false;
        for (u32 v = 0; v < n && !side_entry; ++v) {
            if (!region[v]) continue;
            for (const u32 p : pred[v]) {
                if (p != b && !region[p] && !back.contains({p, v})) {
                    side_entry = true;
                    break;
                }
            }
        }
        shape.side_entries += side_entry;
    }
    return shape;
}

int DumpIr(const fs::path& file) {
    Profile profile = CaptureProfile(file);
    LoadedCapture cap;
    if (const char* error = cap.Load(file)) {
        std::printf("%s: %s\n", file.filename().string().c_str(), error);
        return 2;
    }
    ShaderParams params = cap.Params();
    Info info(HwStage(cap.header.hw_stage), SwStage(cap.header.sw_stage), params);
    Pools pools;
    const auto program = TranslateProgram(cap.code, pools, info, cap.Runtime(), profile);
    for (const IR::Block* block : program.blocks) {
        std::printf("%s\n", IR::DumpBlock(*block).c_str());
    }
    return 0;
}


constexpr float RampLow = -10.0f, RampStops = 24.0f;
constexpr u32 RampWidth = 4096;

std::vector<u32> CompileComposite(const fs::path& file, bool extend, bool ramp,
                                  std::vector<Native::ToneCurveMatch>& matches, std::string& note) {
    Profile profile = CaptureProfile(file);
    LoadedCapture cap;
    if (const char* error = cap.Load(file)) {
        note = error;
        return {};
    }
    ShaderParams params = cap.Params();
    Info info(HwStage(cap.header.hw_stage), SwStage(cap.header.sw_stage), params);
    Pools pools;
    std::vector<IR::Inst*> channels;
    Native::tone_curve_lab_override = extend ? 1 : 0;
    Native::tone_curve_lab_channels = &channels;
    auto program = TranslateProgram(cap.code, pools, info, cap.Runtime(), profile);
    Native::tone_curve_lab_override = -1;
    Native::tone_curve_lab_channels = nullptr;
    matches = Native::FindToneCurves(program);
    if (ramp) {
        if (!Native::IsComposite(matches)) {
            note = "not a composite";
            return {};
        }
        IR::Block& block = *matches[0].result->GetParent();
        std::map<const IR::Inst*, size_t> order;
        size_t index = 0;
        for (IR::Inst& inst : block.Instructions()) {
            order[&inst] = index++;
        }
        std::ranges::sort(matches, [&](const auto& a, const auto& b) {
            return order.at(a.result) < order.at(b.result);
        });
        std::vector<IR::Inst*> outputs;
        if (extend) {
            if (channels.size() != 3) {
                note = "the extension did not apply";
                return {};
            }
            outputs = channels;
        } else {
            for (const auto& m : matches) {
                outputs.push_back(m.result);
            }
        }
        for (const auto& m : matches) {
            IR::IREmitter ir{block, IR::Block::InstructionList::s_iterator_to(*m.scene)};
            const IR::F32 x = ir.GetAttribute(IR::Attribute::FragCoord, 0);
            const IR::F32 stops = IR::F32{ir.FPAdd(
                ir.Imm32(RampLow), IR::F32{ir.FPMul(x, ir.Imm32(RampStops / float(RampWidth)))})};
            const IR::Value scene = ir.FPExp2(stops);
            m.scene->SetArg(m.scene->Arg(0).IsImmediate() ? 1 : 0, scene);
        }
        info.loads.Set(IR::Attribute::FragCoord, 0);
        for (IR::Inst& inst : block.Instructions()) {
            if (inst.GetOpcode() != IR::Opcode::SetAttribute ||
                inst.Arg(0).Attribute() != IR::Attribute::RenderTarget0) {
                continue;
            }
            const u32 comp = inst.Arg(2).U32();
            inst.SetArg(1, comp < 3 ? IR::Value{outputs[comp]} : IR::Value{1.0f});
        }
        Optimization::DeadCodeEliminationPass(program);
    }
    Backend::Bindings bindings{};
    return Backend::SPIRV::EmitSPIRV(profile, cap.Runtime(), program, bindings);
}

int CurvesOne(const fs::path& file) {
    std::vector<Native::ToneCurveMatch> matches, unused;
    std::string note;
    const auto game = CompileComposite(file, false, false, matches, note);
    if (game.empty()) {
        std::printf("%s: %s\n", file.filename().string().c_str(), note.c_str());
        return 2;
    }
    const auto extended = CompileComposite(file, true, false, unused, note);
    const bool composite = Native::IsComposite(matches);
    bool varies = false;
    if (!composite && game != extended) {
        varies = CompileComposite(file, false, false, unused, note) != game ||
                 CompileComposite(file, false, false, unused, note) != extended;
    }
    const char* family = matches.empty() ? "-"
                         : matches[0].family == Native::ToneCurveMatch::Family::Exponential
                             ? "exponential"
                             : "reinhard";
    std::printf("CURVES|channels=%zu;composite=%u;family=%s;words=%u/%u;changed=%u;varies=%u\n",
                matches.size(), composite ? 1u : 0u, family,
                matches.empty() ? 0u : matches[0].k_word,
                matches.empty() ? 0u : matches[0].shoulder_word, game != extended ? 1u : 0u,
                varies ? 1u : 0u);
    return 0;
}

int RampModules(const fs::path& file, const fs::path& out) {
    std::vector<Native::ToneCurveMatch> matches, unused;
    std::string note;
    const auto game = CompileComposite(file, false, true, matches, note);
    const auto hdr = game.empty() ? game : CompileComposite(file, true, true, unused, note);
    if (game.empty() || hdr.empty()) {
        std::printf("%s: %s\n", file.filename().string().c_str(), note.c_str());
        return 1;
    }
    fs::create_directories(out);
    const auto write = [&](const char* name, const std::vector<u32>& words) {
        std::ofstream o(out / name, std::ios::binary);
        o.write(reinterpret_cast<const char*>(words.data()), std::streamsize(words.size() * 4));
    };
    write("game.spv", game);
    write("hdr.spv", hdr);
    std::ofstream meta(out / "ramp.txt");
    meta << (matches[0].family == Native::ToneCurveMatch::Family::Exponential ? "exponential"
                                                                             : "reinhard")
         << ' ' << matches[0].k_word << ' ' << matches[0].shoulder_word << ' ' << RampLow << ' '
         << RampStops << ' ' << RampWidth << ' ' << (g_native ? "uniform" : "storage") << '\n';
    std::printf("%s: ramp modules in %s\n", file.filename().string().c_str(), out.string().c_str());
    return 0;
}

int CensusOne(const fs::path& file) {
    Profile profile = CaptureProfile(file);
    LoadedCapture cap;
    if (const char* error = cap.Load(file)) {
        std::printf("%s: %s\n", file.filename().string().c_str(), error);
        return 2;
    }
    ShaderParams params = cap.Params();
    Info info(HwStage(cap.header.hw_stage), SwStage(cap.header.sw_stage), params);
    Pools pools;
    Native::TablePlan plan;
    Native::TablePlan::lab_target = &plan;
    const auto program = TranslateProgram(cap.code, pools, info, cap.Runtime(), profile);
    Native::TablePlan::lab_target = nullptr;
    const FlowShape flow = AnalyzeFlow(program.ins_list);
    u32 ubo = 0, ssbo = 0, texel = 0, written = 0, special = 0;
    for (const auto& buffer : info.buffers) {
        if (buffer.IsSpecial()) {
            ++special;
        } else if (buffer.is_formatted) {
            ++texel;
        } else if (buffer.only_const_reads && !buffer.is_written) {
            ++ubo;
        } else {
            ++ssbo;
        }
        written += buffer.is_written;
    }
    u32 lds = 0;
    for (const auto& inst : program.ins_list) {
        const auto name = magic_enum::enum_name(inst.opcode);
        lds += name.starts_with("DS_READ") || name.starts_with("DS_WRITE");
    }
    u32 storage_images = 0, array_images = 0, depth_images = 0, native_images = 0;
    for (const auto& image : info.images) {
        storage_images += image.is_written || image.is_atomic;
        array_images += image.is_array;
        depth_images += image.is_depth;
        native_images += image.needs_native;
    }
    std::printf("CENSUS|stage=%u;gcn_insts=%zu;blocks=%u;branches=%u;exec_branches=%u;exits=%u;"
                "loops=%u;max_loop_depth=%u;irreducible=%u;multi_exit_loops=%u;side_entries=%u;"
                "const_buffers=%u;storage_buffers=%u;texel_buffers=%u;written_buffers=%u;"
                "special_buffers=%u;images=%zu;storage_images=%u;array_images=%u;depth_images=%u;"
                "native_images=%u;samplers=%zu;fmasks=%zu;flat_dwords=%u;table_roots=%u;tables=%u;"
                "table_depth=%u;dynamic_offsets=%u;plan_failed=%u;lds=%u;discard=%u;gather=%u;"
                "image_query=%u;fetch_shader=%u;lane_id=%u;quad=%u;shuffle=%u;ballot=%u;fp16=%u;"
                "fp64=%u;dma=%u\n",
                cap.header.hw_stage, program.ins_list.size(), flow.blocks, flow.branches,
                flow.exec_branches, flow.exits, flow.loops, flow.max_loop_depth, flow.irreducible,
                flow.multi_exit_loops, flow.side_entries, ubo, ssbo, texel, written, special,
                info.images.size(), storage_images, array_images, depth_images, native_images,
                info.samplers.size(), info.fmasks.size(), info.srt_info.flattened_bufsize_dw,
                plan.roots, plan.tables, plan.depth, plan.dynamic, plan.valid ? 0u : 1u,
                lds ? 1u : 0u, info.has_discard ? 1u : 0u,
                info.has_image_gather ? 1u : 0u, info.has_image_query ? 1u : 0u,
                info.has_fetch_shader ? 1u : 0u, info.uses_lane_id ? 1u : 0u,
                info.uses_group_quad ? 1u : 0u, info.uses_group_shuffle ? 1u : 0u,
                info.uses_group_ballot ? 1u : 0u, info.uses_fp16 ? 1u : 0u,
                info.uses_fp64 ? 1u : 0u, info.uses_dma ? 1u : 0u);
    if (!plan.valid) {
        std::printf("PLAN|%s\n", plan.error.c_str());
    }
    std::map<std::string_view, u32> ops;
    for (const auto& inst : program.ins_list) {
        ++ops[magic_enum::enum_name(inst.opcode)];
    }
    std::printf("GCN|");
    for (const auto& [name, count] : ops) {
        std::printf("%.*s=%u;", int(name.size()), name.data(), count);
    }
    std::printf("\n");
    return 0;
}

#ifdef _WIN32
bool Readable(u64 at, u64 bytes) {
    const u64 end = at + bytes;
    if (end < at) {
        return false;
    }
    while (at < end) {
        MEMORY_BASIC_INFORMATION mbi{};
        if (!VirtualQuery(reinterpret_cast<const void*>(at), &mbi, sizeof(mbi)) ||
            mbi.State != MEM_COMMIT || (mbi.Protect & (PAGE_NOACCESS | PAGE_GUARD)) ||
            !(mbi.Protect & (PAGE_READONLY | PAGE_READWRITE | PAGE_WRITECOPY | PAGE_EXECUTE_READ |
                             PAGE_EXECUTE_READWRITE | PAGE_EXECUTE_WRITECOPY))) {
            return false;
        }
        at = reinterpret_cast<u64>(mbi.BaseAddress) + mbi.RegionSize;
    }
    return true;
}

LONG CALLBACK WalkerFault(EXCEPTION_POINTERS* ep) {
    if (ep->ExceptionRecord->ExceptionCode == EXCEPTION_ACCESS_VIOLATION &&
        Core::Signals::Instance()->DispatchAccessViolation(
            ep, reinterpret_cast<void*>(ep->ExceptionRecord->ExceptionInformation[1]))) {
        return EXCEPTION_CONTINUE_EXECUTION;
    }
    return EXCEPTION_CONTINUE_SEARCH;
}

int TablesOne(const fs::path& file) {
    Profile profile = CaptureProfile(file);
    LoadedCapture cap;
    if (const char* error = cap.Load(file)) {
        std::printf("%s: %s\n", file.filename().string().c_str(), error);
        return 2;
    }
    ShaderParams params = cap.Params();
    Info info(HwStage(cap.header.hw_stage), SwStage(cap.header.sw_stage), params);
    Pools pools;
    Native::TablePlan plan;
    Native::TablePlan::lab_target = &plan;
    (void)TranslateProgram(cap.code, pools, info, cap.Runtime(), profile);
    Native::TablePlan::lab_target = nullptr;
    const auto walker = info.srt_info.walker_func;
    const u32 size = info.srt_info.flattened_bufsize_dw;
    if (!walker || size <= NUM_USER_DATA_REGS) {
        std::printf("TABLES|no tables\n");
        return 0;
    }
    if (!plan.valid) {
        std::printf("TABLES|no plan;%s\n", plan.error.c_str());
        return 0;
    }
    if (plan.flat_dwords != size) {
        std::printf("TABLES|fail;plan has %u flattened dwords, the walker %u\n", plan.flat_dwords, size);
        return 0;
    }
    Shader::RegisterSrtFaultHandler();
    AddVectoredExceptionHandler(1, WalkerFault);
    {
        Serialization::Archive out;
        plan.Serialize(out);
        Serialization::Archive in{out.TakeOff()};
        Native::TablePlan copy;
        const auto same = [](const auto& a, const auto& b) {
            return a.size() == b.size() &&
                   (a.empty() || std::memcmp(a.data(), b.data(), a.size() * sizeof(a[0])) == 0);
        };
        if (!copy.Deserialize(in) || !copy.valid || copy.flat_dwords != plan.flat_dwords ||
            !same(copy.entries, plan.entries) || !same(copy.ops, plan.ops) ||
            !same(copy.table_dwords, plan.table_dwords)) {
            std::printf("TABLES|fail;the plan read back from its cache record differs\n");
            return 0;
        }
        plan = std::move(copy);
    }

    std::string failures;
    u32 checks = 0, walker_faults = 0, refused = 0, unmapped_faults = 0, unmapped_refused = 0;
    u32 replayed = 0;
    const auto compare = [&](const char* what, const std::array<u32, NUM_USER_DATA_REGS>& user_data,
                             const std::vector<u32>* expected, const std::vector<bool>* skip) {
        std::vector<u32> by_walker(size, 0xCDCDCDCDu), by_plan(size, 0xCDCDCDCDu);
        std::memcpy(by_walker.data(), user_data.data(), sizeof(user_data));
        std::memcpy(by_plan.data(), user_data.data(), sizeof(user_data));
        const u32 faults = Shader::Info::srt_walk_faults;
        walker(user_data.data(), by_walker.data());
        walker_faults += Shader::Info::srt_walk_faults - faults;
        refused += plan.Run(user_data.data(), by_plan.data(), Readable);
        ++checks;
        for (u32 d = NUM_USER_DATA_REGS; d < size; ++d) {
            if (by_walker[d] != by_plan[d]) {
                u32 differ = 0;
                for (u32 e = NUM_USER_DATA_REGS; e < size; ++e) {
                    differ += by_walker[e] != by_plan[e];
                }
                char line[160];
                std::snprintf(line, sizeof(line), "%s: dword %u walker %#x plan %#x (%u of %u differ);",
                              what, d, by_walker[d], by_plan[d], differ, size - NUM_USER_DATA_REGS);
                failures += line;
                return;
            }
        }
        if (expected) {
            for (u32 d = NUM_USER_DATA_REGS; d < size && d < expected->size(); ++d) {
                if ((*skip)[d]) {
                    continue;
                }
                replayed += (*expected)[d] != 0;
                if (by_walker[d] != (*expected)[d]) {
                    char line[160];
                    std::snprintf(line, sizeof(line), "%s: dword %u walker %#x captured %#x;", what, d,
                                  by_walker[d], (*expected)[d]);
                    failures += line;
                    return;
                }
            }
        }
    };

    compare("null", {}, nullptr, nullptr);

    static u32* region = nullptr;
    constexpr size_t RegionBytes = (64u << 20) + (1u << 20);
    for (uintptr_t at = 0x20000000; !region && at < 0xF0000000; at += 0x10000000) {
        region = static_cast<u32*>(VirtualAlloc(reinterpret_cast<void*>(at), RegionBytes,
                                                MEM_RESERVE | MEM_COMMIT, PAGE_READWRITE));
    }
    if (!region) {
        std::printf("TABLES|fail;no low memory for synthetic tables\n");
        return 0;
    }
    const u32 base = u32(reinterpret_cast<uintptr_t>(region));
    for (u32 seed = 1; seed <= 2; ++seed) {
        u32 x = seed * 0x9E3779B9u;
        for (size_t i = 0; i < RegionBytes / 4; ++i) {
            x ^= x << 13;
            x ^= x >> 17;
            x ^= x << 5;
            region[i] = base + ((x % 1024u) << 16);
        }
        std::array<u32, NUM_USER_DATA_REGS> user_data{};
        for (u32 i = 0; i < user_data.size(); ++i) {
            user_data[i] = region[i * 977 + seed];
        }
        compare(seed == 1 ? "synthetic 1" : "synthetic 2", user_data, nullptr, nullptr);
    }

    static void* unmapped = VirtualAlloc(nullptr, 1u << 20, MEM_RESERVE, PAGE_NOACCESS);
    if (unmapped) {
        const u64 address = reinterpret_cast<u64>(unmapped);
        std::array<u32, NUM_USER_DATA_REGS> user_data{};
        for (u32 i = 0; i + 1 < user_data.size(); i += 2) {
            user_data[i] = u32(address);
            user_data[i + 1] = u32(address >> 32);
        }
        const u32 faults_before = walker_faults, refused_before = refused;
        compare("unmapped", user_data, nullptr, nullptr);
        unmapped_faults = walker_faults - faults_before;
        unmapped_refused = refused - refused_before;
    }

    if (cap.flat.size() == size) {
        const std::vector<u32>& flat = cap.flat;
        std::vector<bool> is_table(size), skip(size), nonzero_below(size);
        for (u32 i = 0; i < plan.entries.size(); ++i) {
            is_table[plan.entries[i].parent] = true;
        }
        for (u32 i = u32(plan.entries.size()); i-- > 0;) {
            const u32 d = Native::TablePlan::FirstFlat + i;
            const u16 parent = plan.entries[i].parent;
            if (flat[d] != 0 || nonzero_below[d]) {
                nonzero_below[parent] = true;
            }
        }
        std::vector<u32> offsets(plan.entries.size());
        std::map<u16, u32> extent;
        bool too_large = false;
        for (u32 i = 0; i < plan.entries.size(); ++i) {
            const auto& entry = plan.entries[i];
            u32 offset = entry.offset;
            if (entry.expr_len) {
                offset = plan.Offset(entry, flat.data());
            }
            offsets[i] = offset;
            if (offset > (1u << 20)) {
                too_large = true;
                break;
            }
            extent[entry.parent] = std::max(extent[entry.parent], offset + 2);
        }
        if (too_large) {
            failures += "replay skipped: offset beyond 4 MiB;";
        } else {
            std::map<u16, std::vector<u32>> tables;
            for (const auto& [table, dwords] : extent) {
                tables[table].assign(dwords, 0);
            }
            std::array<u32, NUM_USER_DATA_REGS> user_data = cap.user_data;
            const auto pointer_to = [&](u16 table) -> u64 {
                if (!nonzero_below[table]) {
                    return 0;
                }
                return reinterpret_cast<u64>(tables.at(table).data());
            };
            std::map<std::pair<u16, u32>, u32> placed;
            std::set<std::pair<u16, u32>> pointer_slots;
            bool conflict = false;
            for (u32 i = 0; i < plan.entries.size(); ++i) {
                const u32 d = Native::TablePlan::FirstFlat + i;
                const auto& entry = plan.entries[i];
                if (is_table[d]) {
                    auto& table = tables.at(entry.parent);
                    const u64 address = pointer_to(u16(d));
                    table[offsets[i]] = u32(address);
                    table[offsets[i] + 1] = u32(address >> 32);
                    pointer_slots.insert({entry.parent, offsets[i]});
                    pointer_slots.insert({entry.parent, offsets[i] + 1});
                    skip[d] = true;
                }
            }
            for (u32 i = 0; i < plan.entries.size(); ++i) {
                const u32 d = Native::TablePlan::FirstFlat + i;
                const auto& entry = plan.entries[i];
                auto& table = tables.at(entry.parent);
                const u32 offset = offsets[i];
                if (is_table[d]) {
                    continue;
                }
                if (pointer_slots.contains({entry.parent, offset})) {
                    skip[d] = true;
                    continue;
                }
                const auto [it, added] = placed.try_emplace({entry.parent, offset}, flat[d]);
                conflict |= !added && it->second != flat[d];
                table[offset] = it->second;
            }
            for (const auto& [table, dwords] : tables) {
                if (table < Native::TablePlan::FirstFlat) {
                    const u64 address = pointer_to(table);
                    user_data[table] = u32(address);
                    if (table + 1u < user_data.size()) {
                        user_data[table + 1] = u32(address >> 32);
                    }
                }
            }
            if (conflict) {
                failures += "replay: two flattened values at one table place differ;";
            }
            compare("captured", user_data, &flat, &skip);
        }
    }
    const char* mutant = "none";
    for (size_t i = plan.entries.size(); i-- > 0;) {
        if (plan.entries[i].expr_len != 0) {
            continue;
        }
        Native::TablePlan wrong = plan;
        ++wrong.entries[i].offset;
        std::array<u32, NUM_USER_DATA_REGS> user_data{};
        for (u32 r = 0; r < user_data.size(); ++r) {
            user_data[r] = region[r * 977 + 2];
        }
        std::vector<u32> by_walker(size, 0), by_wrong(size, 0);
        std::memcpy(by_walker.data(), user_data.data(), sizeof(user_data));
        std::memcpy(by_wrong.data(), user_data.data(), sizeof(user_data));
        walker(user_data.data(), by_walker.data());
        wrong.Run(user_data.data(), by_wrong.data(), Readable);
        mutant = by_walker == by_wrong ? "missed" : "caught";
        break;
    }
    if (failures.empty()) {
        std::printf("TABLES|pass;checks=%u;walker_faults=%u;plan_refused=%u;unmapped_faults=%u;"
                    "unmapped_refused=%u;replayed=%u;mutant=%s\n",
                    checks, walker_faults, refused, unmapped_faults, unmapped_refused, replayed,
                    mutant);
    } else {
        std::printf("TABLES|fail;%s\n", failures.c_str());
    }
    return 0;
}

int RunNative(const fs::path& self, const fs::path& dir, const fs::path& report, bool tables) {
    std::vector<fs::path> files;
    for (const auto& entry : fs::directory_iterator(dir)) {
        if (entry.path().extension() == ".bbs") {
            files.push_back(entry.path());
        }
    }
    std::sort(files.begin(), files.end());
    std::ofstream csv(report);
    std::map<std::string, std::pair<u64, u32>> totals;
    std::map<std::string, std::pair<u64, u32>> gcn;
    std::map<std::string, u32> results, plan_errors;
    std::map<std::string, u64> counters;
    std::vector<std::string> keys;
    u32 crashed = 0, structured = 0, censused = 0;
    for (size_t i = 0; i < files.size(); ++i) {
        DWORD code = 0;
        std::wstring args = std::wstring(tables ? L"--tables-one" : L"--census-one") + L" \"" +
                            files[i].wstring() + L"\"";
        if (g_native) args += L" --native";
        const std::string out = RunSelf(self, args, code);
        std::string census, ops, result, plan_error;
        for (size_t at = 0; at < out.size();) {
            const size_t end = out.find('\n', at);
            std::string row = out.substr(at, end == std::string::npos ? std::string::npos : end - at);
            if (!row.empty() && row.back() == '\r') row.pop_back();
            if (row.rfind("CENSUS|", 0) == 0) census = row.substr(7);
            if (row.rfind("GCN|", 0) == 0) ops = row.substr(4);
            if (row.rfind("TABLES|", 0) == 0) result = row.substr(7);
            if (row.rfind("PLAN|", 0) == 0) plan_error = row.substr(5);
            if (end == std::string::npos) break;
            at = end + 1;
        }
        const std::string name = files[i].filename().string();
        if (tables) {
            if (result.empty() || code != 0) {
                ++crashed;
                csv << name << ",crashed (exit " << code << ")\n";
                continue;
            }
            csv << name << ',' << result << '\n';
            ++results[result.substr(0, result.find(';'))];
            for (const char* key : {"walker_faults=", "plan_refused=", "unmapped_faults=",
                                    "unmapped_refused=", "replayed="}) {
                if (const size_t at = result.find(key); at != std::string::npos) {
                    counters[key] += std::strtoull(result.c_str() + at + std::strlen(key), nullptr, 10);
                }
            }
            if (result.find("mutant=caught") != std::string::npos) ++counters["wrong plans caught"];
            if (result.find("mutant=missed") != std::string::npos) ++counters["wrong plans missed"];
            if (result.rfind("no plan;", 0) == 0) ++plan_errors[result.substr(8)];
        } else {
            if (census.empty() || code != 0) {
                ++crashed;
                continue;
            }
            ++censused;
            if (!plan_error.empty()) ++plan_errors[plan_error];
            std::map<std::string, u64> row;
            for (size_t at = 0; at < census.size();) {
                const size_t end = census.find(';', at);
                const std::string pair = census.substr(at, end == std::string::npos ? std::string::npos : end - at);
                const size_t eq = pair.find('=');
                if (eq != std::string::npos) {
                    const std::string key = pair.substr(0, eq);
                    const u64 value = std::strtoull(pair.c_str() + eq + 1, nullptr, 10);
                    if (std::find(keys.begin(), keys.end(), key) == keys.end()) keys.push_back(key);
                    row[key] = value;
                    totals[key].first += value;
                    totals[key].second += value != 0;
                }
                if (end == std::string::npos) break;
                at = end + 1;
            }
            structured += row["irreducible"] == 0 && row["multi_exit_loops"] == 0 && row["side_entries"] == 0;
            if (censused == 1) {
                csv << "file";
                for (const auto& key : keys) csv << ',' << key;
                csv << '\n';
            }
            csv << name;
            for (const auto& key : keys) csv << ',' << row[key];
            csv << '\n';
            for (size_t at = 0; at < ops.size();) {
                const size_t end = ops.find(';', at);
                const std::string pair = ops.substr(at, end == std::string::npos ? std::string::npos : end - at);
                const size_t eq = pair.find('=');
                if (eq != std::string::npos) {
                    auto& entry = gcn[pair.substr(0, eq)];
                    entry.first += std::strtoull(pair.c_str() + eq + 1, nullptr, 10);
                    ++entry.second;
                }
                if (end == std::string::npos) break;
                at = end + 1;
            }
        }
        if ((i + 1) % 200 == 0) std::printf("%zu of %zu\n", i + 1, files.size());
    }
    if (tables) {
        std::printf("Table plans: %zu captures, %u crashed\n", files.size(), crashed);
        for (const auto& [result, n] : results) std::printf("  %s: %u\n", result.c_str(), n);
        for (const auto& [key, n] : counters) {
            std::printf("  %s %llu\n", key.c_str(), (unsigned long long)n);
        }
    } else {
        std::ofstream out(fs::path(report).replace_filename(report.stem().string() + "-gcn.csv"));
        out << "opcode,count,captures\n";
        for (const auto& [name, entry] : gcn) out << name << ',' << entry.first << ',' << entry.second << '\n';
        std::printf("Census: %zu captures, %u read, %u crashed the current recompiler\n", files.size(),
                    censused, crashed);
        std::printf("  structured (reducible, one exit per loop, no side entries): %u\n", structured);
        for (const auto& key : keys) {
            std::printf("  %s: %llu in all, nonzero in %u\n", key.c_str(),
                        (unsigned long long)totals[key].first, totals[key].second);
        }
        std::printf("  GCN opcodes used: %zu\n", gcn.size());
    }
    for (const auto& [error, n] : plan_errors) std::printf("  no plan (%s): %u\n", error.c_str(), n);
    std::printf("Per shader: %s\n", report.string().c_str());
    return 0;
}
int RunCurves(const fs::path& self, const fs::path& dir, const fs::path& report) {
    std::vector<fs::path> files;
    for (const auto& entry : fs::directory_iterator(dir)) {
        if (entry.path().extension() == ".bbs") {
            files.push_back(entry.path());
        }
    }
    std::sort(files.begin(), files.end());
    std::ofstream csv(report);
    csv << "file,result\n";
    u32 crashed = 0, composites = 0, failures = 0, with_channels = 0, varying = 0;
    for (const auto& file : files) {
        DWORD code = 0;
        std::wstring args = L"--curves-one \"" + file.wstring() + L"\"";
        if (g_native) {
            args += L" --native";
        }
        const std::string out = RunSelf(self, args, code);
        const size_t at = out.find("CURVES|");
        const std::string name = file.filename().string();
        if (at == std::string::npos || code != 0) {
            ++crashed;
            csv << name << ",crashed (exit " << code << ")\n";
            continue;
        }
        std::string line = out.substr(at + 7, out.find('\n', at) - at - 7);
        if (!line.empty() && line.back() == '\r') {
            line.pop_back();
        }
        csv << name << ',' << line << '\n';
        const bool composite = line.find("composite=1") != std::string::npos;
        const bool changed = line.find("changed=1") != std::string::npos;
        const bool channels = line.rfind("channels=0", 0) != 0;
        with_channels += channels;
        if (composite) {
            ++composites;
            std::printf("  composite %s: %s\n", name.c_str(), line.c_str());
        }
        const bool varies = line.find("varies=1") != std::string::npos;
        if (varies) {
            ++varying;
            std::printf("  varies %s: the compiler gives other SPIR-V from run to run\n", name.c_str());
        }
        if ((composite != changed && !varies) || (!composite && channels)) {
            ++failures;
            std::printf("  FAIL %s: %s\n", name.c_str(), line.c_str());
        }
    }
    std::printf("Tone curves: %zu captures, %u crashed, %u composites recognised, %u shaders with "
                "curve channels, %u failures (%u compile differently from run to run)\n",
                files.size(), crashed, composites, with_channels, failures, varying);
    std::printf("Per shader: %s\n", report.string().c_str());
    return failures ? 1 : 0;
}
#endif

}

int LiftCache(const fs::path& dir, const fs::path& report, const fs::path& spv_dir, int rounds) {
    const auto profile_bytes = ReadAll(dir / "profile.bin");
    Profile profile{};
    if (profile_bytes.size() != sizeof(Profile)) {
        std::fprintf(stderr, "%s: profile of another engine\n", (dir / "profile.bin").string().c_str());
        return 3;
    }
    std::memcpy(&profile, profile_bytes.data(), sizeof(Profile));
    if (!ApplyLabProfile(profile)) {
        return 2;
    }
    if (!spv_dir.empty()) {
        fs::create_directories(spv_dir / "old");
        fs::create_directories(spv_dir / "new");
    }
    std::ofstream csv(report);
    csv << "name,stage,gcn_dwords,old_ns,new_ns,old_words,new_words,blocks,loops,selections,phis,phis_folded,nodes,live,exec_selects,decline\n";
    Pools pools;
    std::map<std::string, u32> declines;
    u32 total = 0, lifted = 0;
    unsigned long long total_old = 0, total_new = 0, words_old = 0, words_new = 0;
    for (const auto& name : SpvNames(dir)) {
        unsigned long long pgm = 0;
        size_t perm = 0;
        if (std::sscanf(name.c_str(), "0x%llx_%zu", &pgm, &perm) != 2) {
            continue;
        }
        StageSpecialization spec;
        size_t meta_perm = ~size_t{};
        const auto meta = ReadAll(MetaPath(dir, pgm, perm));
        LoadedCapture cap;
        if (meta.empty() ||
            !Vulkan::ReadShaderMetaSpec(std::vector<u8>(meta.begin(), meta.end()), spec, meta_perm) ||
            meta_perm != perm || cap.Load(dir / (name + Native::StateExtension)) || cap.code.empty()) {
            continue;
        }
        Gcn::fetch_override = cap.fetch.empty() && cap.header.fetch_sgpr != ~0u && spec.fetch_shader_data
                                  ? &*spec.fetch_shader_data
                                  : nullptr;
        unsigned long long best_old = ~0ull, best_new = ~0ull;
        std::vector<u32> old_spv;
        Lifter::Output lifted_out;
        for (int round = 0; round < rounds; ++round) {
            ShaderParams params = cap.Params();
            Info info(HwStage(cap.header.hw_stage), SwStage(cap.header.sw_stage), params);
            RuntimeInfo runtime_info = cap.Runtime();
            if (const char* planes = std::getenv("BB_LAB_UCP"); planes && info.hw_stage == HwStage::Vertex) {
                runtime_info.hw.vs.user_clip_plane_mask = u32(std::strtoul(planes, nullptr, 0));
            }
            const RuntimeInfo before = runtime_info;
            const auto t0 = std::chrono::steady_clock::now();
            const auto program = TranslateProgram(cap.code, pools, info, runtime_info, profile);
            Backend::Bindings bindings = spec.start;
            old_spv = Backend::SPIRV::EmitSPIRV(profile, runtime_info, program, bindings);
            const auto t1 = std::chrono::steady_clock::now();
            const Lifter::Input input = Lifter::MakeInput(info, before, profile, spec.start, cap.code);
            const auto t2 = std::chrono::steady_clock::now();
            lifted_out = Lifter::Lift(input);
            if (lifted_out.ok() && lifted_out.uses_tone_curve != info.uses_tone_curve) {
                lifted_out.decline = "tone curve: the two compilers disagree";
                lifted_out.spirv.clear();
            }
            const auto t3 = std::chrono::steady_clock::now();
            best_old = std::min<unsigned long long>(best_old, std::chrono::duration_cast<std::chrono::nanoseconds>(t1 - t0).count());
            best_new = std::min<unsigned long long>(best_new, std::chrono::duration_cast<std::chrono::nanoseconds>(t3 - t2).count());
        }
        ++total;
        const auto& st = lifted_out.stats;
        csv << name << ',' << cap.header.hw_stage << ',' << cap.code.size() << ',' << best_old << ','
            << best_new << ',' << old_spv.size() << ',' << lifted_out.spirv.size() << ',' << st.blocks << ','
            << st.loops << ',' << st.selections << ',' << st.phis << ',' << st.phis_folded << ',' << st.nodes
            << ',' << st.nodes_live << ',' << st.selects_for_exec << ',' << lifted_out.decline << ',' << lifted_out.decline_pc << ',' << lifted_out.last_pc << '\n';
        if (lifted_out.ok()) {
            ++lifted;
            total_old += best_old;
            total_new += best_new;
            words_old += old_spv.size();
            words_new += lifted_out.spirv.size();
            if (!spv_dir.empty()) {
                std::ofstream(spv_dir / "old" / (name + ".spv"), std::ios::binary).write(reinterpret_cast<const char*>(old_spv.data()), old_spv.size() * 4);
                std::ofstream(spv_dir / "new" / (name + ".spv"), std::ios::binary).write(reinterpret_cast<const char*>(lifted_out.spirv.data()), lifted_out.spirv.size() * 4);
            }
        } else {
            ++declines[lifted_out.decline];
        }
    }
    std::printf("%u states, %u lifted\n", total, lifted);
    std::printf("over the lifted: recompiler %.1f ms, lifter %.1f ms; SPIR-V %llu words against %llu\n",
                total_old / 1e6, total_new / 1e6, words_old, words_new);
    for (const auto& [why, count] : declines) {
        std::printf("%5u %s\n", count, why.c_str());
    }
    return 0;
}

int LiftFiles(const fs::path& dir, const fs::path& report, const fs::path& spv_dir) {
    if (!spv_dir.empty()) {
        fs::create_directories(spv_dir);
    }
    std::ofstream csv(report);
    csv << "name,stage,dwords,ns,words,buffers,images,samplers,flat,decline,pc\n";
    std::map<std::string, u32> declines;
    u32 total = 0, lifted = 0;
    unsigned long long total_ns = 0, total_words = 0, total_dwords = 0;
    for (const auto& entry : fs::directory_iterator(dir)) {
        const auto ext = entry.path().extension().string();
        if (ext.size() != 4 || ext[0] != '.') {
            continue;
        }
        const auto bytes = ReadAll(entry.path());
        std::vector<u32> code(bytes.size() / 4);
        std::memcpy(code.data(), bytes.data(), code.size() * 4);
        Lifter::Input in;
        in.discover = true;
        in.code = code;
        const auto decoded = Lifter::Decode(code);
        bool has_lds_write = false, has_ring_store = false, exports_pos = false, exports_color = false, has_interp = false;
        for (const auto& inst : decoded.insts) {
            has_lds_write |= inst.enc == Lifter::Enc::Ds && (inst.op == Lifter::Op::DS_WRITE_B32 || inst.op == Lifter::Op::DS_WRITE2_B32);
            has_ring_store |= inst.enc == Lifter::Enc::Mubuf && inst.glc && inst.slc;
            exports_pos |= inst.enc == Lifter::Enc::Exp && inst.target >= 12 && inst.target < 16;
            exports_color |= inst.enc == Lifter::Enc::Exp && inst.target < 10;
            has_interp |= inst.enc == Lifter::Enc::Vintrp;
        }
        const auto kind = ext.substr(1);
        if (kind == "ppo") {
            in.stage = Lifter::Stage::Fragment;
        } else if (kind == "cpo") {
            in.stage = Lifter::Stage::Compute;
        } else if (kind == "hpo") {
            in.stage = Lifter::Stage::Hull;
        } else if (kind == "dpo") {
            in.stage = Lifter::Stage::Domain;
        } else if (kind == "gpo") {
            in.stage = Lifter::Stage::Geometry;
        } else if (kind == "vpo") {
            in.stage = has_lds_write ? Lifter::Stage::Local : has_ring_store ? Lifter::Stage::Export : Lifter::Stage::Vertex;
        } else {
            in.stage = has_lds_write && !exports_color && !exports_pos ? Lifter::Stage::Local
                       : exports_color || has_interp ? Lifter::Stage::Fragment
                       : exports_pos ? Lifter::Stage::Vertex
                                     : Lifter::Stage::Compute;
        }
        in.num_user_sgprs = 16;
        in.num_input_vgprs = 3;
        for (u32 reg = 0; reg < 16; ++reg) {
            in.ud_push[reg] = s32(reg);
        }
        in.opt.khr_barycentric = true;
        in.opt.uniform_dword_layout = true;
        in.opt.szinp_ok = true;
        for (u32 i = 0; i < 32; ++i) {
            in.ps_inputs.push_back({u8(i), false, false, 0});
        }
        in.z_export_format = 9;
        in.mrtz_mask = 15;
        in.workgroup_size = {1, 1, 1};
        in.tess.ls_stride = 64;
        in.tess.hs_cp_stride = 64;
        in.tess.output_vertices = 3;
        in.tess.domain = 1;
        in.gs.in_vertices = 3;
        in.gs.output_vertices = 4;
        in.gs.invocations = 1;
        in.gs.out_primitive = 2;
        const auto t0 = std::chrono::steady_clock::now();
        const auto out = Lifter::Lift(in);
        const auto ns = std::chrono::duration_cast<std::chrono::nanoseconds>(std::chrono::steady_clock::now() - t0).count();
        ++total;
        csv << entry.path().filename().string() << ',' << int(in.stage) << ',' << code.size() << ',' << ns << ',' << out.spirv.size() << ','
            << out.buffers.size() << ',' << out.images.size() << ',' << out.samplers.size() << ',' << out.flat_table.size() << ','
            << out.decline << ',' << out.last_pc << '\n';
        if (out.ok()) {
            ++lifted;
            total_ns += ns;
            total_words += out.spirv.size();
            total_dwords += code.size();
            if (!spv_dir.empty()) {
                std::ofstream(spv_dir / (entry.path().stem().string() + ".spv"), std::ios::binary).write(reinterpret_cast<const char*>(out.spirv.data()), out.spirv.size() * 4);
            }
        } else {
            ++declines[out.decline];
        }
    }
    std::printf("%u programs, %u lifted in %.1f ms, %llu dwords of GCN to %llu words of SPIR-V\n", total, lifted, total_ns / 1e6, total_dwords, total_words);
    for (const auto& [why, count] : declines) {
        std::printf("%5u %s\n", count, why.c_str());
    }
    return 0;
}

int main(int argc, char** argv) {
    std::vector<char*> args;
    for (int i = 0; i < argc; ++i) {
        if (std::strcmp(argv[i], "--native") == 0) {
            g_native = true;
        } else if (std::strcmp(argv[i], "--spv-dir") == 0 && i + 1 < argc) {
            g_spv_dir = argv[++i];
            fs::create_directories(g_spv_dir);
        } else {
            args.push_back(argv[i]);
        }
    }
    argc = int(args.size());
    argv = args.data();
#ifdef _WIN32
    if (argc >= 3 && (std::strcmp(argv[1], "--recompile-batch") == 0 || std::strcmp(argv[1], "--one") == 0 ||
                      std::strcmp(argv[1], "--walk-null") == 0 || std::strcmp(argv[1], "--census-one") == 0 ||
                      std::strcmp(argv[1], "--tables-one") == 0 ||
                      std::strcmp(argv[1], "--curves-one") == 0 || std::strcmp(argv[1], "--fix-keys") == 0)) {
        SetErrorMode(SEM_FAILCRITICALERRORS | SEM_NOGPFAULTERRORBOX | SEM_NOOPENFILEERRORBOX);
        SetUnhandledExceptionFilter([](EXCEPTION_POINTERS*) -> LONG {
            std::fflush(stdout);
            TerminateProcess(GetCurrentProcess(), 0xC0DE);
            return EXCEPTION_EXECUTE_HANDLER;
        });
    }
#endif
    if (argc >= 3 && std::strcmp(argv[1], "--lift-files") == 0) {
        return LiftFiles(argv[2], argc >= 4 ? fs::path(argv[3]) : fs::path("lab-lift-files.csv"), g_spv_dir);
    }
    if (argc >= 3 && std::strcmp(argv[1], "--lift") == 0) {
        return LiftCache(argv[2], argc >= 4 ? fs::path(argv[3]) : fs::path("lab-lift.csv"), g_spv_dir,
                         argc >= 5 ? std::max(1, std::atoi(argv[4])) : 1);
    }
    if (argc >= 3 && std::strcmp(argv[1], "--check-cache") == 0) {
        return CheckCache(argv[2]);
    }
    if (argc >= 3 && std::strcmp(argv[1], "--fix-keys") == 0) {
        return FixKeys(argv[2]);
    }
    if (argc >= 4 && std::strcmp(argv[1], "--export-catalog") == 0) {
        return ExportCatalog(argv[2], argv[3]);
    }
    if (argc >= 6 && std::strcmp(argv[1], "--seed-cache") == 0) {
#ifdef _WIN32
        wchar_t seed_self[MAX_PATH];
        GetModuleFileNameW(nullptr, seed_self, MAX_PATH);
        return SeedCache(seed_self, argv[2], argv[3], argv[4], argv[5]);
#else
        return SeedCache(argv[0], argv[2], argv[3], argv[4], argv[5]);
#endif
    }
    if (argc >= 4 && std::strcmp(argv[1], "--recompile-batch") == 0) {
        return RecompileBatch(argv[2], std::strtoull(argv[3], nullptr, 10),
                              argc >= 5 ? size_t(std::strtoull(argv[4], nullptr, 10)) : SIZE_MAX,
                              argc >= 6 ? fs::path(argv[5]) : fs::path{});
    }
    if (argc >= 3 && std::strcmp(argv[1], "--bench") == 0) {
#ifdef _WIN32
        _putenv("BB_FRAME_STATS=1");
#else
        setenv("BB_FRAME_STATS", "1", 1);
#endif
        const fs::path dir = argv[2];
        return Bench(dir, argc >= 4 ? std::max(1, std::atoi(argv[3])) : 3,
                     argc >= 5 ? fs::path(argv[4]) : fs::path("lab-bench.csv"));
    }
    if (argc >= 3 && std::strcmp(argv[1], "--recompile-cache") == 0) {
#ifdef _WIN32
        wchar_t self_path[MAX_PATH];
        GetModuleFileNameW(nullptr, self_path, MAX_PATH);
        return RecompileCache(self_path, argv[2]);
#else
        return RecompileCache(argv[0], argv[2]);
#endif
    }
#ifdef _WIN32
    if (argc >= 3 && std::strcmp(argv[1], "--curves-one") == 0) {
        return CurvesOne(argv[2]);
    }
    if (argc >= 4 && std::strcmp(argv[1], "--ramp-modules") == 0) {
        return RampModules(argv[2], argv[3]);
    }
    if (argc >= 3 && std::strcmp(argv[1], "--ir") == 0) {
        return DumpIr(argv[2]);
    }
    if (argc >= 3 && std::strcmp(argv[1], "--census-one") == 0) {
        return CensusOne(argv[2]);
    }
    if (argc >= 3 && std::strcmp(argv[1], "--tables-one") == 0) {
        return TablesOne(argv[2]);
    }
    if (argc >= 3 && std::strcmp(argv[1], "--curves") == 0) {
        wchar_t curves_self[MAX_PATH];
        GetModuleFileNameW(nullptr, curves_self, MAX_PATH);
        const fs::path dir = argv[2];
        return RunCurves(curves_self, dir, argc >= 4 ? fs::path(argv[3]) : dir / "lab-curves.csv");
    }
    if (argc >= 3 && (std::strcmp(argv[1], "--census") == 0 || std::strcmp(argv[1], "--tables") == 0)) {
        wchar_t native_self[MAX_PATH];
        GetModuleFileNameW(nullptr, native_self, MAX_PATH);
        const fs::path dir = argv[2];
        const bool tables = std::strcmp(argv[1], "--tables") == 0;
        const fs::path report = argc >= 4 ? fs::path(argv[3]) : dir / (tables ? "lab-tables.csv" : "lab-census.csv");
        return RunNative(native_self, dir, report, tables);
    }
#endif
    if (argc >= 3 && std::strcmp(argv[1], "--walk-null") == 0) {
        return WalkNull(argv[2]);
    }
    if (argc >= 3 && std::strcmp(argv[1], "--one") == 0) {
        const fs::path file = argv[2];
        return RunOne(file, file.parent_path() / "profile.bin");
    }
    if (argc < 2) {
        std::fprintf(stderr, "usage: bb-shader-lab <capture folder> [report.csv]\n");
        return 2;
    }
    const fs::path dir = argv[1];
    const fs::path report = argc >= 3 ? fs::path(argv[2]) : dir / "lab-report.csv";
#ifdef _WIN32
    wchar_t self[MAX_PATH];
    GetModuleFileNameW(nullptr, self, MAX_PATH);
    return RunAll(self, dir, report);
#else
    return RunAll(argv[0], dir, report);
#endif
}