// igneum-worker-cuda: the one-click NVIDIA worker for igneum-miner --worker. 4 October 2026.
//
// Nothing to install but the NVIDIA driver. The pack's kernels (kernel.cu: cache fill and dataset build;
// kernel_bound.cu: the header-bound hash) are compiled at run time by NVRTC, the toolkit's runtime compiler, which
// ships next to this exe as nvrtc64_120_0.dll plus nvrtc-builtins64_128.dll (NVIDIA's redistributable, see
// THIRD-PARTY.md). The GPU is driven through the driver API in nvcuda.dll, which every NVIDIA driver installs. Both
// libraries are loaded with LoadLibrary/GetProcAddress (cuda_api.h), so no import library is linked and the exe is
// cross-compiled on the Mac with mingw (build-windows.sh). Plain C++17 otherwise.
//
// The source handed to NVRTC is the pack's own text: kernel.cu and kernel_bound.cu up to the host-side launch
// wrappers (which nvcc compiles for the host and NVRTC has no use for), with the pack's program.h and memhard.h as
// named headers, byte for byte. Two stub headers stand in for <cuda_runtime.h> and <cstdint>, which nvcc takes from
// the toolkit. emu/test.sh checks the equality on the Mac. The memory-hard core is therefore the same text host.cu
// compiles, and the miner's CPU re-check of every found nonce covers the rest.
//
// Worker protocol (the same lines as proto-cuda/host.cu --serve and proto-opencl/host.c --serve):
//   stdin:  job <job_id> <header_prehash_hex 64> <target_hex 16> <nonce_start u64> <nonce_count u64> <epoch_seed_hex 64> <day_seed_hex>
//           prepare <epoch_seed_hex 64> <day_seed_hex> <pack_dir>     compile that pack in the background, build its cache
//                                                                     and dataset, self-test it; a job on it then switches
//           quit
//   stdout: ready cuda <device> pack <seed string> dataset-log2 N batch B regs R prepare 1 path nvrtc ...
//           found <job_id> <nonce u64> <hash_hex 16>
//           done <job_id> <hashes> <ms>
//           error <job_id> <text>
//           need <epoch_seed_hex> <day_seed_hex>    before the mismatch error: the pair this worker lacks (the miner prepares it)
//           prepared <epoch_seed_hex> <day_seed_hex> <ms> ...   |   prepare-failed <epoch_seed_hex> <day_seed_hex> <text>
//           info ...
// The first pack comes from --pack <dir> (igneum-miner export-pack writes it; the launcher passes it). Every pack is
// self-tested before it serves a job: cache head, last line and FNV-1a 64, dataset head, last word and 64 samples,
// and the three vector warps of vectors.h through the bound kernel with the pack's own seed words. A pack that fails
// is refused.
//
// Variant racing (4 October 2026, evening; docs/design/miner-tuning.md): every pack's bound kernel is compiled in
// several variants (loop unrolling, the dataset load path: plain, __ldg, __ldcg, __ldcs; a register budget through
// -maxrregcount or __launch_bounds__; threads per block), each self-tested against the pack's vectors (bit-exact or
// discarded) and run for about two seconds on the card; the fastest serves the hour. A race runs inside the prepare
// (the hourly compile-ahead, one lead before the boundary) and never delays the swap: it has a time budget, "base"
// (the pack's text as shipped) is always the first entry, and a prepare that runs out of budget keeps the best so far.
// While a variant is timed the job loop pauses (one mutex): the numbers are exclusive, mining resumes between
// variants. One line per race: `race <epoch16> device <name> ... variants N a=MH/s b=MH/s ... winner <name> <MH/s>
// gain <pct> ...`. A tuning file (--tuning, or IGNEUM_TUNING_FILE from the app) may pin a variant for this card
// model or order the candidates; the app's over-the-air manifest carries it (fleet learning). Under IGNEUM_EMU the
// race is off (the stand-in checks that the handed-over text equals the pack's).
//
// Usage: igneum-worker-cuda --serve --pack <dir> [--device D] [--batch-log2 22] [--block-warps 1] [--arch sm_120|auto]
//                           [--race on|off|<name,name,...>] [--race-bench-ms 2000] [--race-budget-s 120] [--race-rounds 1]
//                           [--variant <name>] [--tuning <file>]
//        igneum-worker-cuda --check --pack <dir> [--device D]      compile, build, self-test, print timings, exit 0/1
//        igneum-worker-cuda --race --pack <dir> [--device D] [--race-rounds 3]   the race alone: one line per variant, exit 0/1

#include <cstdint>
#include <cstdarg>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <chrono>
#include <string>
#include <vector>
#include <thread>
#include <atomic>
#include <mutex>
#include <algorithm>
#include <iostream>

#include "cuda_api.h"
#include "packfile.h"

#ifdef _WIN32
#define WIN32_LEAN_AND_MEAN
#include <windows.h>
#else
#include <dlfcn.h>
#include <dirent.h>
#endif

static const char* WORKER_VERSION = "1.0 (4 October 2026)";

// ---------------------------------------------------------------------------------------------
// Helpers

static double wallMs() {
    using namespace std::chrono;
    return duration<double, std::milli>(steady_clock::now().time_since_epoch()).count();
}

static void emit(const std::string& s) { std::fputs(s.c_str(), stdout); std::fputc('\n', stdout); std::fflush(stdout); }
static void info(const std::string& s) { emit("info " + s); }

static std::string fmt(const char* f, ...) {
    char buf[2048];
    va_list ap;
    va_start(ap, f);
    vsnprintf(buf, sizeof(buf), f, ap);
    va_end(ap);
    return buf;
}

static std::string readText(const std::string& path, bool& ok) {
    size_t n = 0;
    char* b = pf_read_file(path.c_str(), &n);
    if (!b) { ok = false; return ""; }
    std::string s(b, n);
    free(b);
    ok = true;
    return s;
}

static std::string exeDir() {
#ifdef _WIN32
    char buf[MAX_PATH];
    DWORD n = GetModuleFileNameA(nullptr, buf, MAX_PATH);
    std::string p(buf, n);
    size_t i = p.find_last_of("\\/");
    return i == std::string::npos ? "." : p.substr(0, i);
#else
    return ".";
#endif
}

// ---------------------------------------------------------------------------------------------
// Loading the two libraries

static void* libOpen(const std::string& name) {
#ifdef _WIN32
    return (void*)LoadLibraryA(name.c_str());
#else
    return dlopen(name.c_str(), RTLD_NOW);
#endif
}
static void* libSym(void* lib, const char* name) {
#ifdef _WIN32
    return (void*)GetProcAddress((HMODULE)lib, name);
#else
    return dlsym(lib, name);
#endif
}

#define LOAD_SYM(table, field, name) do { table.field = (decltype(table.field))libSym(lib, name); if (!table.field) { missing += std::string(missing.empty() ? "" : ", ") + name; } } while (0)

static bool loadDriver(Drv& d, std::string& err, std::string& libName) {
#ifdef IGNEUM_EMU
    emu_fill_driver(d); libName = "emulation (host threads, no GPU)"; (void)err; return true;
#else
#ifdef _WIN32
    const char* names[] = { "nvcuda.dll" };
#else
    const char* names[] = { "libcuda.so.1", "libcuda.so" };
#endif
    void* lib = nullptr;
    for (const char* n : names) { lib = libOpen(n); if (lib) { libName = n; break; } }
    if (!lib) { err = "the CUDA driver library (nvcuda.dll) is not installed: install or update the NVIDIA driver"; return false; }
    std::string missing;
    LOAD_SYM(d, init, "cuInit");
    LOAD_SYM(d, driverGetVersion, "cuDriverGetVersion");
    LOAD_SYM(d, deviceGetCount, "cuDeviceGetCount");
    LOAD_SYM(d, deviceGet, "cuDeviceGet");
    LOAD_SYM(d, deviceGetName, "cuDeviceGetName");
    LOAD_SYM(d, deviceGetAttribute, "cuDeviceGetAttribute");
    LOAD_SYM(d, deviceTotalMem, "cuDeviceTotalMem_v2");
    LOAD_SYM(d, primaryCtxSetFlags, "cuDevicePrimaryCtxSetFlags_v2");
    LOAD_SYM(d, primaryCtxRetain, "cuDevicePrimaryCtxRetain");
    LOAD_SYM(d, primaryCtxRelease, "cuDevicePrimaryCtxRelease_v2");
    LOAD_SYM(d, ctxSetCurrent, "cuCtxSetCurrent");
    LOAD_SYM(d, ctxSynchronize, "cuCtxSynchronize");
    LOAD_SYM(d, memGetInfo, "cuMemGetInfo_v2");
    LOAD_SYM(d, memAlloc, "cuMemAlloc_v2");
    LOAD_SYM(d, memFree, "cuMemFree_v2");
    LOAD_SYM(d, memcpyDtoH, "cuMemcpyDtoH_v2");
    LOAD_SYM(d, memcpyHtoD, "cuMemcpyHtoD_v2");
    LOAD_SYM(d, moduleLoadData, "cuModuleLoadData");
    LOAD_SYM(d, moduleUnload, "cuModuleUnload");
    LOAD_SYM(d, moduleGetFunction, "cuModuleGetFunction");
    LOAD_SYM(d, launchKernel, "cuLaunchKernel");
    LOAD_SYM(d, streamCreate, "cuStreamCreate");
    LOAD_SYM(d, streamSynchronize, "cuStreamSynchronize");
    LOAD_SYM(d, streamDestroy, "cuStreamDestroy_v2");
    LOAD_SYM(d, funcGetAttribute, "cuFuncGetAttribute");
    LOAD_SYM(d, occupancy, "cuOccupancyMaxActiveBlocksPerMultiprocessor");
    LOAD_SYM(d, getErrorString, "cuGetErrorString");
    LOAD_SYM(d, getErrorName, "cuGetErrorName");
    if (!missing.empty()) { err = "the driver library lacks " + missing + " (driver too old; CUDA 11 or newer is needed)"; return false; }
    return true;
#endif
}

#ifdef _WIN32
// nvrtc64_<major>0_0.dll next to the exe (any major), then a toolkit on PATH. IGNEUM_NVRTC_DLL overrides.
static std::vector<std::string> nvrtcCandidates() {
    std::vector<std::string> v;
    if (const char* o = std::getenv("IGNEUM_NVRTC_DLL")) v.push_back(o);
    std::string dir = exeDir();
    WIN32_FIND_DATAA fd;
    HANDLE h = FindFirstFileA((dir + "\\nvrtc64_*_0.dll").c_str(), &fd);
    if (h != INVALID_HANDLE_VALUE) {
        do { std::string n = fd.cFileName; if (n.find(".alt.") == std::string::npos) v.push_back(dir + "\\" + n); } while (FindNextFileA(h, &fd));
        FindClose(h);
    }
    v.push_back("nvrtc64_120_0.dll");
    v.push_back("nvrtc64_130_0.dll");
    if (const char* cp = std::getenv("CUDA_PATH")) { v.push_back(std::string(cp) + "\\bin\\nvrtc64_120_0.dll"); v.push_back(std::string(cp) + "\\bin\\nvrtc64_130_0.dll"); }
    return v;
}
#endif

static bool loadNvrtc(Rtc& r, std::string& err, std::string& libName) {
#ifdef IGNEUM_EMU
    emu_fill_nvrtc(r); libName = "emulation (source recorded and checked, nothing compiled)"; (void)err; return true;
#else
    void* lib = nullptr;
#ifdef _WIN32
    for (const std::string& n : nvrtcCandidates()) { lib = libOpen(n); if (lib) { libName = n; break; } }
    if (!lib) { err = "nvrtc64_120_0.dll (and nvrtc-builtins64_128.dll) must sit next to " + exeDir() + "\\igneum-worker-cuda.exe; they are in the package"; return false; }
#else
    const char* names[] = { "libnvrtc.so.12", "libnvrtc.so" };
    for (const char* n : names) { lib = libOpen(n); if (lib) { libName = n; break; } }
    if (!lib) { err = "libnvrtc.so.12 not found"; return false; }
#endif
    std::string missing;
    LOAD_SYM(r, version, "nvrtcVersion");
    LOAD_SYM(r, createProgram, "nvrtcCreateProgram");
    LOAD_SYM(r, destroyProgram, "nvrtcDestroyProgram");
    LOAD_SYM(r, compileProgram, "nvrtcCompileProgram");
    LOAD_SYM(r, getProgramLogSize, "nvrtcGetProgramLogSize");
    LOAD_SYM(r, getProgramLog, "nvrtcGetProgramLog");
    LOAD_SYM(r, getPTXSize, "nvrtcGetPTXSize");
    LOAD_SYM(r, getPTX, "nvrtcGetPTX");
    LOAD_SYM(r, getCUBINSize, "nvrtcGetCUBINSize");
    LOAD_SYM(r, getCUBIN, "nvrtcGetCUBIN");
    LOAD_SYM(r, addNameExpression, "nvrtcAddNameExpression");
    LOAD_SYM(r, getLoweredName, "nvrtcGetLoweredName");
    LOAD_SYM(r, getErrorString, "nvrtcGetErrorString");
    if (!missing.empty()) { err = "the NVRTC library lacks " + missing; return false; }
    r.getNumSupportedArchs = (decltype(r.getNumSupportedArchs))libSym(lib, "nvrtcGetNumSupportedArchs");
    r.getSupportedArchs = (decltype(r.getSupportedArchs))libSym(lib, "nvrtcGetSupportedArchs");
    return true;
#endif
}

// ---------------------------------------------------------------------------------------------
// The device context

struct Ctx {
    Drv drv;
    Rtc rtc;
    // read-width experiment, variant 5 (5 October 2026): persistent warps and their scratch; --warps caps the launch
    int warps = 0;                      // 0 = the resident capacity from the occupancy query, rounded down to a power of two
    uint32_t salt = 1;                  // the running per-unit tag salt (+= units per launch)
    int batches = 5;                    // --bench: timed dispatches
    CUdevice dev = 0;
    CUcontext ctx = nullptr;
    std::string name;
    int major = 0, minor = 0, sms = 0, driverVersion = 0, rtcMajor = 0, rtcMinor = 0;
    std::vector<int> rtcArchs;          // what NVRTC can target (empty when the query is unavailable)
    std::string archOpt;                // "sm_120" or "compute_120": what the packs are compiled for
    bool ptx = false;                   // true when archOpt is compute_XY (PTX, driver JIT)
    std::string why;                    // how archOpt was chosen
    int blockWarps = 1;
    // variant racing (see the header): which variants, how long each is timed, the budget of a race, rounds
    std::string race = "on";            // on | off | comma list of variant names
    int raceBenchMs = 2000, raceBudgetS = 120, raceRounds = 1, batchLog2 = 22;
    std::string pinned;                 // --variant: use this variant, no race
    std::string tuning;                 // the tuning file's text ("" = none)
    std::string err(CUresult r) { const char* s = nullptr; if (drv.getErrorString) drv.getErrorString(r, &s); return s ? s : "CUDA driver error"; }
};

#define DRV_CHECK(c, call, what) do { CUresult r_ = (call); if (r_ != CUDA_SUCCESS) { err = std::string(what) + ": " + (c).err(r_); return false; } } while (0)

static bool openDevice(Ctx& c, int device, const std::string& archArg, std::string& err) {
    DRV_CHECK(c, c.drv.init(0), "cuInit");
    int count = 0;
    DRV_CHECK(c, c.drv.deviceGetCount(&count), "cuDeviceGetCount");
    if (count == 0) { err = "no CUDA device"; return false; }
    if (device < 0 || device >= count) { err = fmt("device %d out of range (%d devices)", device, count); return false; }
    DRV_CHECK(c, c.drv.deviceGet(&c.dev, device), "cuDeviceGet");
    char name[256] = {0};
    DRV_CHECK(c, c.drv.deviceGetName(name, 255, c.dev), "cuDeviceGetName");
    c.name = name;
    for (char& ch : c.name) if (ch == ' ') ch = '_';
    DRV_CHECK(c, c.drv.deviceGetAttribute(&c.major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, c.dev), "compute capability major");
    DRV_CHECK(c, c.drv.deviceGetAttribute(&c.minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, c.dev), "compute capability minor");
    DRV_CHECK(c, c.drv.deviceGetAttribute(&c.sms, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT, c.dev), "multiprocessor count");
    c.drv.driverGetVersion(&c.driverVersion);
    // Blocking sync, set before the context exists: the host thread sleeps in cuStreamSynchronize instead of spinning
    // (one full core per worker at the default spin schedule, measured on the RTX 5090 with eight workers, 3 Oct 2026).
    c.drv.primaryCtxSetFlags(c.dev, CU_CTX_SCHED_BLOCKING_SYNC);
    DRV_CHECK(c, c.drv.primaryCtxRetain(&c.ctx, c.dev), "cuDevicePrimaryCtxRetain");
    DRV_CHECK(c, c.drv.ctxSetCurrent(c.ctx), "cuCtxSetCurrent");
    c.rtc.version(&c.rtcMajor, &c.rtcMinor);
    if (c.rtc.getNumSupportedArchs && c.rtc.getSupportedArchs) {
        int n = 0;
        if (c.rtc.getNumSupportedArchs(&n) == NVRTC_SUCCESS && n > 0 && n < 256) { c.rtcArchs.assign((size_t)n, 0); if (c.rtc.getSupportedArchs(c.rtcArchs.data()) != NVRTC_SUCCESS) c.rtcArchs.clear(); }
    }
    // The target: the device's own SASS (sm_XY) when this NVRTC knows the architecture, else PTX for the newest
    // architecture it knows below the device's, which the driver JIT-compiles forward. --arch overrides.
    int cc = c.major * 10 + c.minor;
    if (archArg != "auto" && !archArg.empty()) {
        c.archOpt = archArg; c.ptx = archArg.rfind("compute_", 0) == 0; c.why = "--arch";
    } else if (c.rtcArchs.empty()) {
        c.archOpt = fmt("sm_%d", cc); c.why = "the device's architecture (NVRTC did not list its targets)";
    } else {
        bool known = false; int best = 0;
        for (int a : c.rtcArchs) { if (a == cc) known = true; if (a <= cc && a > best) best = a; }
        if (known) { c.archOpt = fmt("sm_%d", cc); c.why = "the device's architecture, listed by NVRTC"; }
        else if (best > 0) { c.archOpt = fmt("compute_%d", best); c.ptx = true; c.why = fmt("this NVRTC does not know sm_%d; PTX for compute_%d, JIT-compiled by the driver", cc, best); }
        else { c.archOpt = fmt("compute_%d", c.rtcArchs.front()); c.ptx = true; c.why = fmt("this NVRTC knows nothing at or below sm_%d; PTX for its oldest target", cc); }
    }
    return true;
}

// ---------------------------------------------------------------------------------------------
// NVRTC: compile one of the pack's kernel files

// The pack's kernel files end in host-side launch wrappers (cudaError_t igneum_launch_* with <<< >>> launches) that
// nvcc compiles for the host. NVRTC compiles device code only, so the text is cut there. The cut is checked: nothing
// device-side may follow it.
static bool deviceOnly(const std::string& text, std::string& out, std::string& err) {
    size_t cut = text.find("\n// Host-side launch wrappers");
    if (cut == std::string::npos) cut = text.find("\ncudaError_t ");
    if (cut == std::string::npos) { out = text; return true; }
    std::string tail = text.substr(cut + 1);
    if (tail.find("__global__") != std::string::npos || tail.find("__device__") != std::string::npos) { err = "device code after the host launch wrappers; the pack layout is not the one this worker knows"; return false; }
    out = text.substr(0, cut + 1);
    return true;
}

static const char* STUB_CUDA_RUNTIME =
    "// igneum-worker-cuda: stand-in for <cuda_runtime.h> under NVRTC, which has the device built-ins already\n"
    "#pragma once\n"
    "#ifndef __CUDACC_RTC__\n#error \"this stub is for NVRTC only\"\n#endif\n"
    "#ifdef __SIZE_TYPE__\ntypedef __SIZE_TYPE__ size_t;\n#elif defined(__LP64__) || defined(_LP64)\ntypedef unsigned long size_t;\n#else\ntypedef unsigned long long size_t;\n#endif\n";
static const char* STUB_CSTDINT =
    "// igneum-worker-cuda: stand-in for <cstdint> under NVRTC (the fixed-width types the packs use)\n"
    "#pragma once\n"
    "typedef signed char int8_t; typedef unsigned char uint8_t; typedef short int16_t; typedef unsigned short uint16_t;\n"
    "typedef int int32_t; typedef unsigned int uint32_t;\n"
    "#if defined(__LP64__) || defined(_LP64)\ntypedef long int64_t; typedef unsigned long uint64_t;\n"
    "#else\ntypedef long long int64_t; typedef unsigned long long uint64_t;\n#endif\n";

struct Compiled {
    std::vector<char> image;
    std::vector<std::string> lowered;
    double ms = 0;
    std::string log;
};

static bool rtcCompile(Ctx& c, const std::string& src, const char* name, const std::string& programH, const std::string& memhardH,
                       const std::vector<std::string>& nameExprs, Compiled& out, std::string& err, const std::vector<std::string>& extraOpts = {}) {
    double t0 = wallMs();
    const char* headers[4] = { STUB_CUDA_RUNTIME, STUB_CSTDINT, programH.c_str(), memhardH.c_str() };
    const char* names[4] = { "cuda_runtime.h", "cstdint", "program.h", "memhard.h" };
    nvrtcProgram prog = nullptr;
    nvrtcResult r = c.rtc.createProgram(&prog, src.c_str(), name, 4, headers, names);
    if (r != NVRTC_SUCCESS) { err = std::string("nvrtcCreateProgram: ") + c.rtc.getErrorString(r); return false; }
    for (const std::string& e : nameExprs) {
        r = c.rtc.addNameExpression(prog, e.c_str());
        if (r != NVRTC_SUCCESS) { err = "nvrtcAddNameExpression " + e + ": " + c.rtc.getErrorString(r); c.rtc.destroyProgram(&prog); return false; }
    }
    std::string archOpt = "--gpu-architecture=" + c.archOpt;
    // -default-device: NVRTC rejects unannotated functions as host code (nvcc treats them as host and discards them);
    // the pack headers (program.h, memhard.h) carry plain inline helpers, so every unannotated function is device code here.
    std::vector<const char*> opts = { archOpt.c_str(), "--std=c++17", "-default-device" };
    for (const std::string& o : extraOpts) opts.push_back(o.c_str());
    r = c.rtc.compileProgram(prog, (int)opts.size(), opts.data());
    {
        size_t logSize = 0;
        if (c.rtc.getProgramLogSize(prog, &logSize) == NVRTC_SUCCESS && logSize > 1) {
            std::vector<char> log(logSize);
            c.rtc.getProgramLog(prog, log.data());
            out.log.assign(log.data(), logSize - 1);
        }
    }
    if (r != NVRTC_SUCCESS) {
        std::string one;
        for (char ch : out.log) { if (ch == '\n' || ch == '\r') { if (one.size() && one.back() != '|') one += " | "; } else one += ch; if (one.size() > 600) break; }
        err = std::string("nvrtcCompileProgram ") + name + " for " + c.archOpt + ": " + c.rtc.getErrorString(r) + ": " + one;
        c.rtc.destroyProgram(&prog);
        return false;
    }
    for (const std::string& e : nameExprs) {
        const char* lowered = nullptr;
        r = c.rtc.getLoweredName(prog, e.c_str(), &lowered);
        if (r != NVRTC_SUCCESS || !lowered) { err = "nvrtcGetLoweredName " + e + ": " + c.rtc.getErrorString(r); c.rtc.destroyProgram(&prog); return false; }
        out.lowered.push_back(lowered);
    }
    size_t n = 0;
    if (c.ptx) {
        r = c.rtc.getPTXSize(prog, &n);
        if (r == NVRTC_SUCCESS) { out.image.resize(n); r = c.rtc.getPTX(prog, out.image.data()); }
    } else {
        r = c.rtc.getCUBINSize(prog, &n);
        if (r == NVRTC_SUCCESS) { out.image.resize(n); r = c.rtc.getCUBIN(prog, out.image.data()); }
    }
    c.rtc.destroyProgram(&prog);
    if (r != NVRTC_SUCCESS || n == 0) { err = std::string(c.ptx ? "nvrtcGetPTX" : "nvrtcGetCUBIN") + ": " + c.rtc.getErrorString(r); return false; }
    out.ms = wallMs() - t0;
    return true;
}

// ---------------------------------------------------------------------------------------------
// A resident pair: one pack compiled, its cache and dataset on the device, self-tested

static bool hexEq(const std::string& a, const std::string& b) {
    if (a.size() != b.size()) return false;
    for (size_t i = 0; i < a.size(); ++i) if (std::tolower((unsigned char)a[i]) != std::tolower((unsigned char)b[i])) return false;
    return true;
}

struct Pair;
// A pair is the pair of a job when the job's seeds (the hex the node sent) are the pair's seeds. The derived seed
// words are no identity: a retried program's words are its attempt's words, not the bare seed's (packfile.h,
// 5 October 2026), so comparing words refused every job of a retried program.
static bool pairIs(const Pair* p, const std::string& epochHex, const std::string& dayHex);

struct Pair {
    std::string dir, epochHex, dayHex, seedString;
    uint32_t sw[8] = {0}, kw[8] = {0};
    uint32_t datasetLog2 = 0, words = 0, cacheWords = 0, cacheSegments = 0;
    uint32_t items = 0, mulshift = 0;   // research class ds55: words / 16 items; 1 when the pack's loads are (src * words) >> 32
    CUmodule modKernel = nullptr, modBound = nullptr;
    CUfunction fCacheFill = nullptr, fBuild = nullptr, fHashBound = nullptr;
    CUdeviceptr cache = 0, ds = 0;
    double compileMs = 0, cacheMs = 0, dsMs = 0, checkMs = 0;
    std::string check;
    bool checkPass = false, checked = false;
    int regs = 0, blocksPerSM = 0;
    int blockWarps = 1;                 // threads per block = 32 x this (the winning variant's, else the worker's default)
    std::string variant = "base";       // the bound kernel in service: a variant name (see allVariants)
    std::string raceLine;               // the race's one-line report, emitted by the main thread with "prepared"
    double raceMs = 0;
    // read-width experiment (5 October 2026): the pack's load class and, for variant 5, the persistent-warp scratch
    std::string loadClass = "v2";
    std::string programClass = "v2", eraHex;   // Counter ASIC 2.0: the pack's class and era seed (packfile.h)
    uint32_t loadsPerHash = 128, bytesPerHash = 512, scratchOps = 0;
    bool persistent = false;
    CUdeviceptr scratch = 0;
    int warps = 0;                      // persistent warps launched (the arena holds this many)
    int residentWarps = 0;              // the occupancy query's capacity: blocks/SM x warps/block x SMs
    size_t scratchBytes = 0;
    // hot-table experiment (5 October 2026, docs/plans/hot-table.md): the epoch's hot table, filled on the device by the
    // pack's igneum_hot_fill, the argument after the init words
    uint32_t hotMb = 0, hotWords = 0, hotSegments = 0, hotSlots = 0;
    CUfunction fHotFill = nullptr;
    // class v5 (7 October 2026): the pack's state leaves, uploaded for igneum_build and freed after it (0 for other classes)
    uint32_t stateLeaves = 0;
    CUdeviceptr hot = 0;
    double hotMs = 0;
};

static bool pairIs(const Pair* p, const std::string& epochHex, const std::string& dayHex) {
    return p && hexEq(p->epochHex, epochHex) && hexEq(p->dayHex, dayHex);
}

// Counter ASIC 2.0: a job that names a class (and an era) belongs to a pair of that class (and era) only, so a pack
// of the old class for the same seeds is not this job's pair and the prepared pack of the right class wins.
static bool pairIsClass(const Pair* p, const std::string& epochHex, const std::string& dayHex, const std::string& cls, const std::string& era) {
    char why[256];
    return pairIs(p, epochHex, dayHex) && pf_pack_class_ok(p->programClass.c_str(), p->eraHex.c_str(), cls.c_str(), era.c_str(), why, sizeof(why));
}

// The trailing `class=` and `era=` tokens of a job or prepare line (absent on every class v2 line), removed from `f`.
static void takeClassTokens(std::vector<std::string>& f, std::string& cls, std::string& era) {
    while (!f.empty()) {
        char c[8] = {0}, e[65] = {0};
        if (!pf_class_token(f.back().c_str(), c, sizeof(c), e, sizeof(e))) break;
        if (c[0]) cls = c;
        if (e[0]) era = e;
        f.pop_back();
    }
}

// The job loop and a race take turns on the card: a variant is timed with no job running (exclusive numbers), and
// mining resumes between variants. Held per chunk by the job loop, per variant by the race.
static std::mutex gpuMutex;

static void releasePair(Ctx& c, Pair* p) {
    if (!p) return;
    if (p->ds) c.drv.memFree(p->ds);
    if (p->cache) c.drv.memFree(p->cache);
    if (p->scratch) c.drv.memFree(p->scratch);
    if (p->hot) c.drv.memFree(p->hot);
    if (p->modBound) c.drv.moduleUnload(p->modBound);
    if (p->modKernel) c.drv.moduleUnload(p->modKernel);
    delete p;
}

struct IgneumInitWordsArg { uint32_t w[8]; };

// `block` threads per block (32 x warps); `nonces` must be a multiple of it.
static bool launchHash(Ctx& c, Pair* p, CUdeviceptr out, uint32_t baseNonce, const uint32_t iw[8], uint32_t nonces, uint32_t block, CUstream s, std::string& err) {
    uint32_t mask = p->words - 1u;
    IgneumInitWordsArg a; std::memcpy(a.w, iw, 32);
    if (p->persistent) {
        // Variant 5: N persistent warps over nonces / 32 units; the arena was sized for p->warps warps in buildPair.
        uint32_t units = nonces / 32u, warps = (uint32_t)p->warps;
        if (warps > units) warps = units;
        while (warps > 1u && units % warps != 0u) warps >>= 1;
        if (block != 32u) { err = "a variant-5 pack runs one warp per block (--block-warps 1)"; return false; }
        uint32_t salt = c.salt; c.salt += units;
        // the hot table (when the pack has one) sits between the init words and the scratch triple
        void* args[9] = { &p->ds, &out, &baseNonce, &mask, &a, &p->scratch, &units, &salt, nullptr };
        if (p->hot) { args[5] = &p->hot; args[6] = &p->scratch; args[7] = &units; args[8] = &salt; }
        DRV_CHECK(c, c.drv.launchKernel(p->fHashBound, warps, 1, 1, 32, 1, 1, 0, s, args, nullptr), "cuLaunchKernel igneum_hash_bound (persistent)");
        return true;
    }
    void* args[6] = { &p->ds, &out, &baseNonce, &mask, &a, &p->hot };
    DRV_CHECK(c, c.drv.launchKernel(p->fHashBound, nonces / block, 1, 1, block, 1, 1, 0, s, args, nullptr), "cuLaunchKernel igneum_hash_bound");
    return true;
}


// ---------------------------------------------------------------------------------------------
// Variant racing

struct Variant {
    std::string name;
    int unroll = 0;      // 0: the iteration loop as emitted; N: "#pragma unroll N" before it (8 = fully unrolled)
    int load = 0;        // 0: plain ds[i]; 1: __ldg (read-only data path); 2: __ldcg (L2 only, no L1); 3: __ldcs (streaming)
    int maxrreg = 0;     // 0: none; N: --maxrregcount=N (registers per thread, occupancy against spills)
    int blockWarps = 0;  // 0: the worker's --block-warps; N: 32 x N threads per block
    int minBlocks = 0;   // N > 0: __launch_bounds__(32 x blockWarps, N) (the compiler fits N blocks per SM)
};

// The catalogue. Names are stable: the tuning file and the fleet records use them. "base" is the pack's text as
// shipped with the worker's default block and is always the first entry of a race.
static std::vector<Variant> allVariants() {
    std::vector<Variant> v;
    auto add = [&](const char* n, int unroll, int load, int maxrreg, int bw, int minBlocks) { Variant x; x.name = n; x.unroll = unroll; x.load = load; x.maxrreg = maxrreg; x.blockWarps = bw; x.minBlocks = minBlocks; v.push_back(x); };
    add("base", 0, 0, 0, 0, 0);
    add("w2", 0, 0, 0, 2, 0);
    add("w4", 0, 0, 0, 4, 0);
    add("w8", 0, 0, 0, 8, 0);
    add("u2", 2, 0, 0, 0, 0);
    add("u8", 8, 0, 0, 0, 0);
    add("ldg", 0, 1, 0, 0, 0);
    add("ldcg", 0, 2, 0, 0, 0);
    add("ldcs", 0, 3, 0, 0, 0);
    add("r32", 0, 0, 32, 0, 0);
    add("r64", 0, 0, 64, 0, 0);
    add("lb4-w4", 0, 0, 0, 4, 4);
    add("lb8-w2", 0, 0, 0, 2, 8);
    add("u2-ldg", 2, 1, 0, 0, 0);
    add("u2-w4", 2, 0, 0, 4, 0);
    add("ldg-w4", 0, 1, 0, 4, 0);
    add("ldcg-w4", 0, 2, 0, 4, 0);
    return v;
}

static const Variant* findVariant(const std::vector<Variant>& all, const std::string& name) {
    for (const Variant& v : all) if (v.name == name) return &v;
    return nullptr;
}

// The variant's source: the pack's bound-kernel text with the variant's rewrites. Every rewrite has an exact anchor
// in the text igneum-pow emits; a text without the anchor refuses the variant (why), it is never guessed.
static bool variantSource(const std::string& base, const Variant& v, int blockWarps, std::string& out, std::string& why) {
    out = base;
    if (v.unroll > 0) {
        const char* anchor = "\n    for (uint32_t it = 0u; it < ";
        size_t p = out.find(anchor);
        if (p == std::string::npos) { why = "no iteration loop in the bound kernel text"; return false; }
        out.insert(p + 1, fmt("#pragma unroll %d\n", v.unroll));
    }
    if (v.load > 0) {
        const char* fn = v.load == 1 ? "__ldg" : v.load == 2 ? "__ldcg" : "__ldcs";
        size_t body = out.find("igneum_hash_bound(");
        if (body == std::string::npos) { why = "no igneum_hash_bound in the text"; return false; }
        size_t p = body; int n = 0;
        while ((p = out.find(" ^ ds[", p)) != std::string::npos) {
            size_t close = out.find(']', p);
            if (close == std::string::npos) { why = "an unterminated dataset load"; return false; }
            out.insert(close + 1, ")");                                 // " ^ ds[idx]"  ->  " ^ __ldg(&ds[idx])"
            out.insert(p + 3, std::string(fn) + "(&");
            p += 6; ++n;
        }
        if (n == 0) { why = "no dataset loads in the bound kernel"; return false; }
    }
    if (v.minBlocks > 0) {
        const char* a = "__global__ void igneum_hash_bound(";
        size_t p = out.find(a);
        if (p == std::string::npos) { why = "no kernel declaration anchor"; return false; }
        out.replace(p, std::strlen(a), fmt("__global__ void __launch_bounds__(%d, %d) igneum_hash_bound(", 32 * blockWarps, v.minBlocks));
    }
    return true;
}

// The tuning file: {"cards": {"<device name as this worker prints it>": {"variant": "u2-ldg", "race": false,
// "candidates": ["u2-ldg", "ldg", "base"]}}, ...}. Read with plain string scanning (no JSON library in this exe);
// a file that does not parse means no tuning. Keys and names are [A-Za-z0-9_.-].
struct Tuning {
    bool found = false;
    std::string variant;                // pinned variant ("" = none)
    bool race = true;                   // false: use the pinned variant without a race
    std::vector<std::string> candidates;
};

static std::string jsonStringAfter(const std::string& t, size_t from, const char* key, size_t limit) {
    size_t k = t.find(std::string("\"") + key + "\"", from);
    if (k == std::string::npos || k > limit) return "";
    size_t q = t.find('"', t.find(':', k) + 1);
    if (q == std::string::npos) return "";
    size_t e = t.find('"', q + 1);
    return e == std::string::npos ? "" : t.substr(q + 1, e - q - 1);
}

static Tuning readTuning(const std::string& text, const std::string& device) {
    Tuning tu;
    if (text.empty()) return tu;
    size_t cards = text.find("\"cards\"");
    if (cards == std::string::npos) return tu;
    size_t k = text.find("\"" + device + "\"", cards);
    if (k == std::string::npos) return tu;
    size_t open = text.find('{', k);
    if (open == std::string::npos) return tu;
    size_t close = open; int depth = 0;
    for (; close < text.size(); ++close) { if (text[close] == '{') ++depth; else if (text[close] == '}' && --depth == 0) break; }
    if (close >= text.size()) return tu;
    tu.found = true;
    tu.variant = jsonStringAfter(text, open, "variant", close);
    size_t r = text.find("\"race\"", open);
    if (r != std::string::npos && r < close) { size_t c = text.find(':', r); tu.race = text.compare(text.find_first_not_of(" \t\r\n", c + 1), 5, "false") != 0; }
    size_t cand = text.find("\"candidates\"", open);
    if (cand != std::string::npos && cand < close) {
        size_t a = text.find('[', cand), b = text.find(']', a == std::string::npos ? cand : a);
        if (a != std::string::npos && b != std::string::npos && b < close) {
            size_t i = a;
            while ((i = text.find('"', i + 1)) != std::string::npos && i < b) { size_t e = text.find('"', i + 1); if (e == std::string::npos || e > b) break; tu.candidates.push_back(text.substr(i + 1, e - i - 1)); i = e; }
        }
    }
    return tu;
}

struct RaceEntry {
    Variant v;
    int blockWarps = 1;                 // the block this entry runs with
    Compiled cb;
    CUmodule mod = nullptr;
    CUfunction fn = nullptr;
    int regs = 0, blocksPerSM = 0;
    double mhs = 0;                     // best round
    bool ok = false;                    // compiled, loaded, self-tested
    std::string note;                   // why not, or a detail
    std::string src;
};

// One timed window on the card for an entry: the pack's vector warps (bit-exact or the entry is out), then
// launches of `batch` nonces until benchMs elapsed (the first launch warms up and is not counted). Holds gpuMutex.
static bool raceTime(Ctx& c, Pair* p, const PfPack& pk, RaceEntry& e, CUdeviceptr dOut, uint32_t batch, int benchMs, CUstream s, bool selfTest) {
    std::lock_guard<std::mutex> hold(gpuMutex);
    CUfunction keep = p->fHashBound;
    p->fHashBound = e.fn;
    std::string err;
    uint32_t block = 32u * (uint32_t)e.blockWarps;
    bool ok = true;
    if (selfTest && pk.haveVectors) {
        std::vector<uint64_t> vec(32);
        for (int w = 0; w < pk.vecWarps && ok; ++w) {
            if (!launchHash(c, p, dOut, pk.vecBase[w], p->sw, block, block, s, err)) { e.note = "launch: " + err; ok = false; break; }
            CUresult r = c.drv.streamSynchronize(s);
            if (r == CUDA_SUCCESS) r = c.drv.memcpyDtoH(vec.data(), dOut, 32u * 8u);
            if (r != CUDA_SUCCESS) { e.note = "vector warp: " + c.err(r); ok = false; break; }
            for (int l = 0; l < 32; ++l) if (vec[(size_t)l] != pk.vecOut[w][l]) { e.note = fmt("vector warp %d lane %d: device %016llx expected %016llx (discarded)", w, l, (unsigned long long)vec[(size_t)l], (unsigned long long)pk.vecOut[w][l]); ok = false; break; }
        }
    }
    if (ok) {
        uint32_t iw[8]; std::memcpy(iw, p->sw, 32);
        uint32_t n = batch - (batch % block);
        if (n == 0) n = block;
        double t0 = 0; uint64_t hashes = 0; int launches = 0;
        while (true) {
            if (!launchHash(c, p, dOut, 0x10000000u + (uint32_t)launches * n, iw, n, block, s, err)) { e.note = "launch: " + err; ok = false; break; }
            CUresult r = c.drv.streamSynchronize(s);
            if (r != CUDA_SUCCESS) { e.note = "bench: " + c.err(r); ok = false; break; }
            double now = wallMs();
            if (launches == 0) t0 = now; else hashes += n;
            ++launches;
            if (launches >= 3 && now - t0 >= benchMs) { double mhs = (double)hashes / (now - t0) / 1000.0; if (mhs > e.mhs) e.mhs = mhs; break; }
        }
    }
    p->fHashBound = keep;
    return ok;
}

// Races the bound kernel of `p` (its cache and dataset are built, its base kernel self-tested) and installs the
// winner: p->modBound, fHashBound, regs, blockWarps, variant. The pair keeps serving its base kernel if every other
// entry fails. `boundDev`, `programH`, `memhardH` are the texts the base was compiled from. Sets p->raceLine.
static void racePair(Ctx& c, Pair* p, const PfPack& pk, const std::string& boundDev, const std::string& programH, const std::string& memhardH, CUstream s) {
    double t0 = wallMs();
    std::vector<Variant> all = allVariants();
    Tuning tu = readTuning(c.tuning, c.name);
    std::string pinned = !c.pinned.empty() ? c.pinned : (tu.found && !tu.race ? tu.variant : "");
    // the order: base first, then the pinned or tuned candidates, then the rest (or the --race list only)
    std::vector<Variant> order;
    auto push = [&](const std::string& n) { const Variant* v = findVariant(all, n); if (v && !findVariant(order, n)) order.push_back(*v); };
    push("base");
    if (!pinned.empty()) push(pinned);
    else {
        for (const std::string& n : tu.candidates) push(n);
        if (c.race != "on" && c.race != "off") { std::string rest = c.race; size_t i = 0; while (i <= rest.size()) { size_t j = rest.find(',', i); if (j == std::string::npos) j = rest.size(); if (j > i) push(rest.substr(i, j - i)); i = j + 1; } }
        else if (c.race == "on") for (const Variant& v : all) push(v.name);
    }
#ifdef IGNEUM_EMU
    order.resize(1);   // the stand-in checks that the handed-over text is the pack's; no rewrites under emulation
#endif
    if (order.size() < 2) {
        // nothing to race against base (emulation, or --race with no known name): no timing, the base kernel serves
        p->blockWarps = c.blockWarps; p->variant = "base"; p->raceMs = wallMs() - t0;
        p->raceLine = fmt("race %.16s device %s variants 1 base only, no race (%s)", p->epochHex.c_str(), c.name.c_str(),
#ifdef IGNEUM_EMU
                          "emulation");
#else
                          "no other variant named");
#endif
        return;
    }
    const bool pinnedOnly = !pinned.empty() && order.size() == 2;
    uint32_t batch = 1u << c.batchLog2;
    int benchMs = c.raceBenchMs;
    double deadline = t0 + c.raceBudgetS * 1000.0;
    std::vector<RaceEntry> entries;
    for (const Variant& v : order) {
        RaceEntry e; e.v = v; e.blockWarps = v.blockWarps > 0 ? v.blockWarps : c.blockWarps;
        if (v.name == "base") { e.mod = p->modBound; e.fn = p->fHashBound; e.regs = p->regs; e.blocksPerSM = p->blocksPerSM; e.ok = true; }
        else if (!variantSource(boundDev, v, e.blockWarps, e.src, e.note)) e.ok = false;
        else e.ok = true;   // compiled below
        entries.push_back(std::move(e));
    }
    // Compile the variants, up to four at a time (NVRTC is thread-safe; the compile is CPU work)
    {
        std::vector<size_t> todo;
        for (size_t i = 1; i < entries.size(); ++i) if (entries[i].ok) todo.push_back(i);
        size_t next = 0;
        std::mutex m;
        auto work = [&]() {
            while (true) {
                size_t i;
                { std::lock_guard<std::mutex> g(m); if (next >= todo.size() || wallMs() > deadline - benchMs) return; i = todo[next++]; }
                RaceEntry& e = entries[i];
                std::vector<std::string> extra;
                if (e.v.maxrreg > 0) extra.push_back(fmt("--maxrregcount=%d", e.v.maxrreg));
                std::string err;
                if (!rtcCompile(c, e.src, "kernel_bound.cu", programH, memhardH, { "igneum_hash_bound" }, e.cb, err, extra)) { e.ok = false; e.note = "compile: " + err.substr(0, 200); }
            }
        };
        int threads = (int)std::min<size_t>(4, std::max<size_t>(1, todo.size()));
        std::vector<std::thread> ts;
        for (int t = 0; t < threads; ++t) ts.emplace_back(work);
        for (std::thread& t : ts) t.join();
        for (size_t i = 1; i < entries.size(); ++i) if (entries[i].ok && entries[i].cb.image.empty()) { entries[i].ok = false; entries[i].note = "not compiled: the race budget ran out"; }
    }
    // Load the modules (the context is current on this thread)
    for (size_t i = 1; i < entries.size(); ++i) {
        RaceEntry& e = entries[i];
        if (!e.ok) continue;
        CUresult r = c.drv.moduleLoadData(&e.mod, e.cb.image.data());
        if (r != CUDA_SUCCESS) { e.ok = false; e.note = "cuModuleLoadData: " + c.err(r); e.mod = nullptr; continue; }
        if (c.drv.moduleGetFunction(&e.fn, e.mod, e.cb.lowered[0].c_str()) != CUDA_SUCCESS) { e.ok = false; e.note = "function not in the module"; continue; }
        c.drv.funcGetAttribute(&e.regs, CU_FUNC_ATTRIBUTE_NUM_REGS, e.fn);
        c.drv.occupancy(&e.blocksPerSM, e.fn, 32 * e.blockWarps, 0);
    }
    double compileMs = wallMs() - t0;
    // Time them: rounds over the entries, interleaved, best per entry. A pinned variant is only self-tested.
    CUdeviceptr dOut = 0;
    std::string benchErr;
    if (c.drv.memAlloc(&dOut, (size_t)batch * 8u) != CUDA_SUCCESS) { benchErr = "cuMemAlloc for the race"; for (RaceEntry& e : entries) if (e.v.name != "base") e.ok = false; }
    int rounds = pinnedOnly ? 1 : std::max(1, c.raceRounds);
    for (int round = 0; round < rounds && benchErr.empty(); ++round) {
        for (size_t i = 0; i < entries.size(); ++i) {
            RaceEntry& e = entries[i];
            if (!e.ok) continue;
            if (i > 0 && round == 0 && wallMs() > deadline) { e.ok = false; e.note = "not timed: the race budget ran out"; continue; }
            if (pinnedOnly && i == 0) continue;
            if (!raceTime(c, p, pk, e, dOut, batch, pinnedOnly ? 0 : benchMs, s, round == 0)) e.ok = false;
            // the mutex is not fair: give the job loop the card between windows (measured on the Mac, 4 October
            // 2026: without this a queued job waited the whole race, 36 s)
            std::this_thread::sleep_for(std::chrono::milliseconds(150));
        }
    }
    if (dOut) c.drv.memFree(dOut);
    // The winner: the fastest entry; base keeps its place unless a variant is at least 0.5% faster (noise guard).
    size_t win = 0;
    if (pinnedOnly && entries.size() == 2 && entries[1].ok) win = 1;
    else for (size_t i = 1; i < entries.size(); ++i) if (entries[i].ok && entries[i].mhs > entries[win].mhs * (win == 0 ? 1.005 : 1.0)) win = i;
    double baseMhs = entries[0].mhs, winMhs = entries[win].mhs;
    if (win != 0) {
        RaceEntry& w = entries[win];
        c.drv.moduleUnload(p->modBound);
        p->modBound = w.mod; p->fHashBound = w.fn; p->regs = w.regs; p->blocksPerSM = w.blocksPerSM; p->blockWarps = w.blockWarps; p->variant = w.v.name;
        w.mod = nullptr;
    } else {
        p->blockWarps = c.blockWarps; p->variant = "base";
    }
    for (size_t i = 1; i < entries.size(); ++i) if (entries[i].mod) c.drv.moduleUnload(entries[i].mod);
    p->raceMs = wallMs() - t0;
    // The one line. Variants in race order: name=MH/s (regs), or name=- (why).
    uint32_t loads = 0, wide = 0;
    pf_define_u32(programH.c_str(), "IGNEUM_LOADS_PER_HASH", &loads);
    pf_define_u32(programH.c_str(), "IGNEUM_WIDE_LOADS_PER_HASH", &wide);
    std::string line = fmt("race %.16s device %s driver %d.%d arch %s loads %u wide %u variants %zu", p->epochHex.c_str(), c.name.c_str(), c.driverVersion / 1000, (c.driverVersion % 100) / 10, c.archOpt.c_str(), loads, wide, entries.size());
    for (const RaceEntry& e : entries) {
        if (e.ok && (e.mhs > 0 || pinnedOnly)) line += fmt(" %s=%.3f/%dr/%dw", e.v.name.c_str(), e.mhs, e.regs, e.blockWarps);
        else line += fmt(" %s=-", e.v.name.c_str());
    }
    line += fmt(" winner %s %.3f base %.3f gain %+.2f%% compile %.0f bench %.0f total %.0f ms%s%s", p->variant.c_str(), winMhs, baseMhs, baseMhs > 0 ? (winMhs / baseMhs - 1.0) * 100.0 : 0.0, compileMs, p->raceMs - compileMs, p->raceMs,
                 pinnedOnly ? " pinned by tuning" : (tu.found ? " tuned order" : ""), benchErr.empty() ? "" : (" " + benchErr).c_str());
    for (const RaceEntry& e : entries) if (!e.ok && !e.note.empty()) line += " | " + e.v.name + ": " + e.note;
    p->raceLine = line;
}

// Compiles the pack in `dir`, builds its cache and dataset on stream `s`, runs the self-test, races the variants
// (`race`). Returns the pair or null with `err` set. Runs on the main thread for --pack and --check, on the prepare
// thread for `prepare`.
static Pair* buildPair(Ctx& c, const std::string& dir, CUstream s, std::string& err, bool race) {
    PfPack pk;
    char perr[512];
    if (!pf_load(dir.c_str(), &pk, perr, sizeof(perr))) { err = std::string("pack ") + dir + ": " + perr; return nullptr; }
    bool ok1, ok2, ok3, ok4;
    std::string kernelCu = readText(dir + "/kernel.cu", ok1), boundCu = readText(dir + "/kernel_bound.cu", ok2);
    std::string programH = readText(dir + "/program.h", ok3), memhardH = readText(dir + "/memhard.h", ok4);
    if (!ok1 || !ok2 || !ok3 || !ok4) { err = "pack " + dir + " lacks kernel.cu, kernel_bound.cu, program.h or memhard.h"; return nullptr; }
    std::string kernelDev, boundDev;
    if (!deviceOnly(kernelCu, kernelDev, err) || !deviceOnly(boundCu, boundDev, err)) { err = "pack " + dir + ": " + err; return nullptr; }
    Pair* p = new Pair();
    p->dir = dir; p->epochHex = pk.epochHex; p->dayHex = pk.dayHex; p->seedString = pk.seedString;
    std::memcpy(p->sw, pk.seedw, 32); std::memcpy(p->kw, pk.keyw, 32);
    // the dataset's size is the pack's word count (IGNEUM_DATASET_WORDS when present, else 1 << IGNEUM_DATASET_LOG2: packfile.h),
    // never the power of two of the floor: a 5.5 GiB research pack holds 1,476,395,008 words and 92,274,688 items
    p->datasetLog2 = pk.datasetLog2; p->words = pk.datasetWords; p->items = pk.datasetItems; p->mulshift = pk.datasetMulshift;
    p->cacheWords = 1u << pk.cacheLog2Words; p->cacheSegments = pk.cacheSegments;
    // Compile
    Compiled ck, cb;
    std::vector<std::string> kernelNames = { "igneum_cache_fill", "igneum_build" };
    if (pk.hotMb) kernelNames.push_back("igneum_hot_fill");
    if (!rtcCompile(c, kernelDev, "kernel.cu", programH, memhardH, kernelNames, ck, err)) { releasePair(c, p); return nullptr; }
    if (!rtcCompile(c, boundDev, "kernel_bound.cu", programH, memhardH, { "igneum_hash_bound" }, cb, err)) { releasePair(c, p); return nullptr; }
    p->compileMs = ck.ms + cb.ms;
    // Load
    {
        CUresult r = c.drv.moduleLoadData(&p->modKernel, ck.image.data());
        if (r != CUDA_SUCCESS) { err = "cuModuleLoadData kernel.cu (" + c.archOpt + "): " + c.err(r); releasePair(c, p); return nullptr; }
        r = c.drv.moduleLoadData(&p->modBound, cb.image.data());
        if (r != CUDA_SUCCESS) { err = "cuModuleLoadData kernel_bound.cu (" + c.archOpt + "): " + c.err(r); releasePair(c, p); return nullptr; }
        if (c.drv.moduleGetFunction(&p->fCacheFill, p->modKernel, ck.lowered[0].c_str()) != CUDA_SUCCESS) { err = "igneum_cache_fill (" + ck.lowered[0] + ") not in the module"; releasePair(c, p); return nullptr; }
        if (c.drv.moduleGetFunction(&p->fBuild, p->modKernel, ck.lowered[1].c_str()) != CUDA_SUCCESS) { err = "igneum_build (" + ck.lowered[1] + ") not in the module"; releasePair(c, p); return nullptr; }
        if (c.drv.moduleGetFunction(&p->fHashBound, p->modBound, cb.lowered[0].c_str()) != CUDA_SUCCESS) { err = "igneum_hash_bound (" + cb.lowered[0] + ") not in the module"; releasePair(c, p); return nullptr; }
        if (pk.hotMb && c.drv.moduleGetFunction(&p->fHotFill, p->modKernel, ck.lowered[2].c_str()) != CUDA_SUCCESS) { err = "igneum_hot_fill (" + ck.lowered[2] + ") not in the module"; releasePair(c, p); return nullptr; }
        c.drv.funcGetAttribute(&p->regs, CU_FUNC_ATTRIBUTE_NUM_REGS, p->fHashBound);
        c.drv.occupancy(&p->blocksPerSM, p->fHashBound, 32 * c.blockWarps, 0);
    }
    p->loadClass = pk.loadClass; p->loadsPerHash = pk.loadsPerHash; p->bytesPerHash = pk.bytesPerHash; p->scratchOps = pk.scratchOps;
    p->programClass = pk.programClass; p->eraHex = pk.eraHex;
    p->stateLeaves = pk.stateLeaves;
    p->persistent = pk.persistent != 0;
    // Class v5 (docs/design/class-v5-stored-state.md): the window's leaves (leaves.bin), read and checked against the pack's
    // count and FNV-1a 64 before anything is allocated; uploaded for igneum_build below and freed right after it, so device
    // memory while hashing is the class v4 worker's. A v5 pack without its leaves builds nothing (the known-failed case).
    uint32_t* hLeaves = nullptr;
    size_t leavesBytes = 0;
    { char lerr[700]; if (!pf_load_leaves(dir.c_str(), &pk, &hLeaves, &leavesBytes, lerr, sizeof(lerr))) { err = "pack " + dir + ": " + lerr; releasePair(c, p); return nullptr; } }
    p->hotMb = pk.hotMb; p->hotWords = pk.hotWords; p->hotSegments = pk.hotSegments; p->hotSlots = pk.hotSlots;
    p->residentWarps = p->blocksPerSM * c.blockWarps * c.sms;
    size_t scratchBytes = 0, hotBytes = (size_t)pk.hotWords * 4u;
    if (p->persistent) {
        // Variant 5: one arena per launched warp. The launch is the resident capacity (the occupancy query), rounded
        // down to a power of two so it divides every batch, or --warps; the allocation cannot change the occupancy
        // (registers and shared memory decide it), and the number is re-queried after the allocation below to show it.
        if (c.blockWarps != 1) { err = "a variant-5 pack runs one warp per block: use --block-warps 1"; releasePair(c, p); return nullptr; }
        int w = c.warps > 0 ? c.warps : p->residentWarps;
        int pw = 1; while (pw * 2 <= w) pw *= 2;
        p->warps = pw;
        scratchBytes = (size_t)p->warps * 32u * (size_t)pk.scratchWordsPerLane * 4u;
    }
    // Cache
    double t0 = wallMs();
    size_t cacheBytes = (size_t)p->cacheWords * 4u, dsBytes = (size_t)p->words * 4u;
    {
        size_t freeB = 0, totalB = 0;
        if (c.drv.memGetInfo(&freeB, &totalB) == CUDA_SUCCESS && freeB < cacheBytes + dsBytes + scratchBytes + hotBytes + leavesBytes + (64u << 20)) {
            err = fmt("%llu MiB free on the device, this pack needs %llu MiB (cache %llu + dataset %llu + scratch %llu + hot %llu + leaves %llu)", (unsigned long long)(freeB >> 20), (unsigned long long)((cacheBytes + dsBytes + scratchBytes + hotBytes + leavesBytes) >> 20), (unsigned long long)(cacheBytes >> 20), (unsigned long long)(dsBytes >> 20), (unsigned long long)(scratchBytes >> 20), (unsigned long long)(hotBytes >> 20), (unsigned long long)(leavesBytes >> 20));
            std::free(hLeaves); releasePair(c, p); return nullptr;
        }
    }
    if (p->persistent) {
        CUresult r = c.drv.memAlloc(&p->scratch, scratchBytes);
        if (r != CUDA_SUCCESS) { err = "cuMemAlloc scratch: " + c.err(r); p->scratch = 0; std::free(hLeaves); releasePair(c, p); return nullptr; }
        p->scratchBytes = scratchBytes;
        int after = 0;
        c.drv.occupancy(&after, p->fHashBound, 32 * c.blockWarps, 0);
        info(fmt("variant 5: %d persistent warps (resident capacity %d = %d blocks/SM x %d warps/block x %d SMs; occupancy query after the allocation %d blocks/SM), scratch %llu MiB (%u KiB per warp)",
                 p->warps, p->residentWarps, p->blocksPerSM, c.blockWarps, c.sms, after, (unsigned long long)(scratchBytes >> 20), pk.scratchWordsPerLane * 4u * 32u / 1024u));
    }
    {
        CUresult r = c.drv.memAlloc(&p->cache, cacheBytes);
        if (r != CUDA_SUCCESS) { err = "cuMemAlloc cache: " + c.err(r); p->cache = 0; std::free(hLeaves); releasePair(c, p); return nullptr; }
        uint32_t nSeg = p->cacheSegments, block = 256u, grid = (nSeg + block - 1u) / block;
        void* args[2] = { &p->cache, &nSeg };
        r = c.drv.launchKernel(p->fCacheFill, grid, 1, 1, block, 1, 1, 0, s, args, nullptr);
        if (r == CUDA_SUCCESS) r = c.drv.streamSynchronize(s);
        if (r != CUDA_SUCCESS) { err = "cache fill: " + c.err(r); std::free(hLeaves); releasePair(c, p); return nullptr; }
    }
    p->cacheMs = wallMs() - t0;
    // Dataset
    t0 = wallMs();
    {
        CUresult r = c.drv.memAlloc(&p->ds, dsBytes);
        if (r != CUDA_SUCCESS) { err = "cuMemAlloc dataset: " + c.err(r); p->ds = 0; std::free(hLeaves); releasePair(c, p); return nullptr; }
        uint32_t nItems = p->items, block = 256u, grid = (nItems + block - 1u) / block;
        if (hLeaves) {
            // class v5: igneum_build(ds, cache, leaves, nLeaves, nItems), the leaf buffer freed once the build has run
            CUdeviceptr dLeaves = 0;
            uint32_t nLeaves = pk.stateLeaves;
            r = c.drv.memAlloc(&dLeaves, leavesBytes);
            if (r != CUDA_SUCCESS) { err = "cuMemAlloc state leaves: " + c.err(r); std::free(hLeaves); releasePair(c, p); return nullptr; }
            r = c.drv.memcpyHtoD(dLeaves, hLeaves, leavesBytes);
            std::free(hLeaves); hLeaves = nullptr;
            if (r == CUDA_SUCCESS) {
                void* args[5] = { &p->ds, &p->cache, &dLeaves, &nLeaves, &nItems };
                r = c.drv.launchKernel(p->fBuild, grid, 1, 1, block, 1, 1, 0, s, args, nullptr);
            }
            if (r == CUDA_SUCCESS) r = c.drv.streamSynchronize(s);
            c.drv.memFree(dLeaves);
            if (r != CUDA_SUCCESS) { err = "dataset build (class v5, " + std::to_string(nLeaves) + " state leaves): " + c.err(r); releasePair(c, p); return nullptr; }
        } else {
            void* args[5] = { &p->ds, &p->cache, &nItems, nullptr, nullptr };   // five slots: the driver reads the kernel's three, the emulation reads the shape from the two null tails
            r = c.drv.launchKernel(p->fBuild, grid, 1, 1, block, 1, 1, 0, s, args, nullptr);
            if (r == CUDA_SUCCESS) r = c.drv.streamSynchronize(s);
            if (r != CUDA_SUCCESS) { err = "dataset build: " + c.err(r); releasePair(c, p); return nullptr; }
        }
    }
    p->dsMs = wallMs() - t0;
    // Hot table (hot-table experiment): filled from the epoch seed by the pack's own kernel, never shipped
    if (pk.hotMb) {
        t0 = wallMs();
        CUresult r = c.drv.memAlloc(&p->hot, hotBytes);
        if (r != CUDA_SUCCESS) { err = "cuMemAlloc hot table: " + c.err(r); p->hot = 0; releasePair(c, p); return nullptr; }
        uint32_t nSeg = pk.hotSegments, block = 256u, grid = (nSeg + block - 1u) / block;
        void* args[2] = { &p->hot, &nSeg };
        r = c.drv.launchKernel(p->fHotFill, grid, 1, 1, block, 1, 1, 0, s, args, nullptr);
        if (r == CUDA_SUCCESS) r = c.drv.streamSynchronize(s);
        if (r != CUDA_SUCCESS) { err = "hot table fill: " + c.err(r); releasePair(c, p); return nullptr; }
        p->hotMs = wallMs() - t0;
    }
    // Self-test against vectors.h
    t0 = wallMs();
    if (!pk.haveVectors) {
        p->checked = false; p->checkPass = true;
        p->check = "self-test skipped (no vectors.h in the pack); the miner's CPU re-check covers every found nonce";
    } else {
        uint32_t cacheHead[16], cacheLast[16], dsHead[16], dsLast = 0;
        std::vector<uint32_t> samples((size_t)(pk.nSamples > 0 ? pk.nSamples : 1), 0u);
        std::vector<uint64_t> vec((size_t)pk.vecWarps * 32u, 0ull);
        std::vector<uint32_t> whole(p->cacheWords);
        CUresult r = c.drv.memcpyDtoH(whole.data(), p->cache, cacheBytes);
        if (r != CUDA_SUCCESS) { err = "cuMemcpyDtoH cache: " + c.err(r); releasePair(c, p); return nullptr; }
        std::memcpy(cacheHead, whole.data(), 64);
        std::memcpy(cacheLast, whole.data() + p->cacheWords - 16u, 64);
        uint64_t fnv = pf_fnv1a64(whole.data(), cacheBytes);
        whole.clear(); whole.shrink_to_fit();
        if ((r = c.drv.memcpyDtoH(dsHead, p->ds, 64)) != CUDA_SUCCESS) { err = "cuMemcpyDtoH dataset head: " + c.err(r); releasePair(c, p); return nullptr; }
        if (pk.dsLastIndex < p->words) r = c.drv.memcpyDtoH(&dsLast, p->ds + (CUdeviceptr)pk.dsLastIndex * 4u, 4);
        for (int i = 0; i < pk.nSamples && r == CUDA_SUCCESS; ++i) if (pk.sampleIdx[i] < p->words) r = c.drv.memcpyDtoH(&samples[(size_t)i], p->ds + (CUdeviceptr)pk.sampleIdx[i] * 4u, 4);
        if (r != CUDA_SUCCESS) { err = "cuMemcpyDtoH dataset words: " + c.err(r); releasePair(c, p); return nullptr; }
        CUdeviceptr out = 0;
        if ((r = c.drv.memAlloc(&out, 32u * (size_t)c.blockWarps * 8u)) != CUDA_SUCCESS) { err = "cuMemAlloc vector out: " + c.err(r); releasePair(c, p); return nullptr; }
        for (int w = 0; w < pk.vecWarps; ++w) {
            // One block of 32 x block-warps lanes; the vector warp is its first 32 lanes (lane nonce = base + gid)
            if (!launchHash(c, p, out, pk.vecBase[w], p->sw, 32u * (uint32_t)c.blockWarps, 32u * (uint32_t)c.blockWarps, s, err)) { c.drv.memFree(out); releasePair(c, p); return nullptr; }
            if ((r = c.drv.streamSynchronize(s)) != CUDA_SUCCESS || (r = c.drv.memcpyDtoH(&vec[(size_t)w * 32u], out, 32u * 8u)) != CUDA_SUCCESS) { err = "vector warp: " + c.err(r); c.drv.memFree(out); releasePair(c, p); return nullptr; }
        }
        c.drv.memFree(out);
        uint32_t hotHead[16] = {0}, hotLast[16] = {0};
        uint64_t hotFnv = 0;
        if (p->hot) {
            std::vector<uint32_t> hw(p->hotWords);
            if ((r = c.drv.memcpyDtoH(hw.data(), p->hot, hotBytes)) != CUDA_SUCCESS) { err = "cuMemcpyDtoH hot table: " + c.err(r); releasePair(c, p); return nullptr; }
            std::memcpy(hotHead, hw.data(), 64);
            std::memcpy(hotLast, hw.data() + p->hotWords - 16u, 64);
            hotFnv = pf_fnv1a64(hw.data(), hotBytes);
        }
        char line[1024];
        p->checkPass = pf_selftest(&pk, cacheHead, cacheLast, fnv, dsHead, dsLast, samples.data(), vec.data(), p->hot ? hotHead : nullptr, p->hot ? hotLast : nullptr, hotFnv, line, sizeof(line)) != 0;
        p->checked = true;
        p->check = line;
    }
    p->checkMs = wallMs() - t0;
    if (!p->checkPass) { err = p->check; releasePair(c, p); return nullptr; }
    p->blockWarps = c.blockWarps;
    if (race && c.race != "off") racePair(c, p, pk, boundDev, programH, memhardH, s);
    return p;
}

static std::string pairSummary(const Pair* p) {
    return fmt("nvrtc %.0f cache %.0f dataset %.0f hot %.0f check %.0f race %.0f ms variant %s class %s%s%s; %s", p->compileMs, p->cacheMs, p->dsMs, p->hotMs, p->checkMs, p->raceMs, p->variant.c_str(), p->programClass.c_str(),
               p->stateLeaves ? fmt(" (state leaves %u, uploaded for the build and freed)", p->stateLeaves).c_str() : "",
               p->mulshift ? fmt(" (dataset %u words, %u items, not a power of two: loads are (src * words) >> 32)", p->words, p->items).c_str() : "", p->check.c_str());
}

// ---------------------------------------------------------------------------------------------
// Prepare, on its own thread

struct PrepareTask {
    std::string epochHex, dayHex, dir, error;
    std::string wantClass, wantEra;     // the class and era the prepare line named (empty: any)
    std::atomic<bool> done{false};
    Pair* result = nullptr;
    double t0 = 0;
    std::thread thread;
};

static void prepareRun(Ctx* c, PrepareTask* t) {
    CUstream s = nullptr;
    std::string err;
    if (c->drv.ctxSetCurrent(c->ctx) != CUDA_SUCCESS) { t->error = "cuCtxSetCurrent on the prepare thread"; t->done = true; return; }
    if (c->drv.streamCreate(&s, CU_STREAM_NON_BLOCKING) != CUDA_SUCCESS) { t->error = "cuStreamCreate on the prepare thread"; t->done = true; return; }
    Pair* p = buildPair(*c, t->dir, s, err, true);
    c->drv.streamDestroy(s);
    if (p && (p->epochHex != t->epochHex || p->dayHex != t->dayHex)) {
        err = "the pack in " + t->dir + " is for epoch " + p->epochHex.substr(0, 16) + " day " + p->dayHex + ", not the prepared seeds";
        releasePair(*c, p); p = nullptr;
    }
    char why[256];
    if (p && !pf_pack_class_ok(p->programClass.c_str(), p->eraHex.c_str(), t->wantClass.c_str(), t->wantEra.c_str(), why, sizeof(why))) {
        err = "pack " + t->dir + ": " + why;
        releasePair(*c, p); p = nullptr;
    }
    t->result = p;
    t->error = err;
    t->done = true;
}

// ---------------------------------------------------------------------------------------------
// Serve

struct Options {
    bool serve = false, check = false, raceOnly = false;
    bool bench = false, memprobe = false;   // read-width experiment (5 October 2026)
    int batches = 5, warps = 0, probeMib = 0;
    int device = 0, batchLog2 = 22, blockWarps = 1;
    std::string pack, arch = "auto";
    std::string race = "on", pinned, tuningPath;
    int raceBenchMs = 2000, raceBudgetS = 120, raceRounds = 0;   // rounds 0 = 1 in --serve, 3 in --race
};

static void usage() {
    std::printf("igneum-worker-cuda %s\n"
                "  --serve --pack <dir>   GPU worker for igneum-miner --worker: jobs on stdin, found/done lines on stdout\n"
                "  --check --pack <dir>   compile the pack, build its cache and dataset, self-test, print timings, exit 0 or 1\n"
                "  --device D             CUDA device index (default 0)\n"
                "  --batch-log2 B         nonces per dispatch = 2^B (default 22)\n"
                "  --block-warps W        warps per thread block (default 1)\n"
                "  --arch sm_XY|compute_XY|auto   NVRTC target (default auto: the device's architecture)\n"
                "  --bench --pack <dir>   read-width experiment: build and self-test the pack, time --batches dispatches of 2^B nonces,\n"
                "                         print the 2^B fingerprint at base nonce 0 (one RESULT line); a variant-5 pack runs --warps persistent warps\n"
                "  --memprobe [--probe-mib N]   no pack: dependent random 4, 16 and 64-byte reads, independent reads, a coalesced stream and an\n"
                "                         integer chain at 4, 64 and 1024 MiB (the same table as igneum-worker-opencl --memprobe)\n"
                "  --batches N            --bench: timed dispatches (default 5)\n"
                "  --warps N              --bench on a variant-5 pack: persistent warps (default: the occupancy capacity, rounded down to a power of two)\n"
                "  --race --pack <dir>    the variant race alone (3 rounds): one line per variant, the race line, exit 0 or 1\n"
                "  --race on|off|a,b,c    in --serve: race every variant (default), none, or these names\n"
                "  --race-bench-ms N      timed window per variant (default 2000)\n"
                "  --race-budget-s N      a race stops compiling and timing after this (default 120; base is kept)\n"
                "  --race-rounds N        interleaved rounds, best per variant (default 1 in --serve, 3 in --race)\n"
                "  --variant <name>       use this variant without a race (also from the tuning file)\n"
                "  --tuning <file>        the per-card tuning file (default: IGNEUM_TUNING_FILE from the environment)\n", WORKER_VERSION);
}

static Options parseArgs(int argc, char** argv) {
    Options o;
    for (int i = 1; i < argc; ++i) {
        std::string a = argv[i];
        auto next = [&]() -> std::string { if (i + 1 >= argc) { usage(); std::exit(2); } return argv[++i]; };
        if (a == "--serve") o.serve = true;
        else if (a == "--check") o.check = true;
        else if (a == "--bench") o.bench = true;
        else if (a == "--memprobe") o.memprobe = true;
        else if (a == "--batches") o.batches = std::atoi(next().c_str());
        else if (a == "--warps") o.warps = std::atoi(next().c_str());
        else if (a == "--probe-mib") o.probeMib = std::atoi(next().c_str());
        else if (a == "--race" && (i + 1 >= argc || std::string(argv[i + 1]).rfind("--", 0) == 0)) o.raceOnly = true;
        else if (a == "--race") o.race = next();
        else if (a == "--race-bench-ms") o.raceBenchMs = std::atoi(next().c_str());
        else if (a == "--race-budget-s") o.raceBudgetS = std::atoi(next().c_str());
        else if (a == "--race-rounds") o.raceRounds = std::atoi(next().c_str());
        else if (a == "--variant") o.pinned = next();
        else if (a == "--tuning") o.tuningPath = next();
        else if (a == "--pack") o.pack = next();
        else if (a == "--device") o.device = std::atoi(next().c_str());
        else if (a == "--batch-log2") o.batchLog2 = std::atoi(next().c_str());
        else if (a == "--block-warps") o.blockWarps = std::atoi(next().c_str());
        else if (a == "--arch") o.arch = next();
        else if (a == "--no-prepare") { /* accepted for symmetry with the other workers; prepare is always on here */ }
        else if (a == "-h" || a == "--help") { usage(); std::exit(0); }
        else { std::printf("unknown argument %s\n", argv[i]); usage(); std::exit(2); }
    }
    if (o.batchLog2 < 10 || o.batchLog2 > 28) { std::printf("--batch-log2 must be between 10 and 28\n"); std::exit(2); }
    if (o.blockWarps < 1 || o.blockWarps > 32) { std::printf("--block-warps must be between 1 and 32\n"); std::exit(2); }
    if (!o.serve && !o.check && !o.raceOnly && !o.bench && !o.memprobe) { usage(); std::exit(2); }
    if (o.raceBenchMs < 200 || o.raceBenchMs > 20000) { std::printf("--race-bench-ms must be between 200 and 20000\n"); std::exit(2); }
    if (o.raceBudgetS < 5 || o.raceBudgetS > 540) { std::printf("--race-budget-s must be between 5 and 540 (the prepare lead is 600 DAA)\n"); std::exit(2); }
    if (o.raceRounds == 0) o.raceRounds = o.raceOnly ? 3 : 1;
    if (o.tuningPath.empty()) if (const char* t = std::getenv("IGNEUM_TUNING_FILE")) o.tuningPath = t;
    if (o.pack.empty() && !o.memprobe) { std::printf("--pack <dir> is required (igneum-miner export-pack <node> <dir> writes one)\n"); std::exit(2); }
    while (o.pack.size() > 1 && (o.pack.back() == '/' || o.pack.back() == '\\')) o.pack.pop_back();
    return o;
}

static std::vector<std::string> split(const std::string& line) {
    std::vector<std::string> f;
    size_t i = 0;
    while (i < line.size()) {
        while (i < line.size() && (line[i] == ' ' || line[i] == '\t' || line[i] == '\r')) ++i;
        size_t j = i;
        while (j < line.size() && line[j] != ' ' && line[j] != '\t' && line[j] != '\r') ++j;
        if (j > i) f.push_back(line.substr(i, j - i));
        i = j;
    }
    return f;
}

// A pack directory for the given seeds under `root` (one subdirectory per pack, each with seeds.txt), or "".
static std::string findPackFor(const std::string& root, const std::string& epochHex, const std::string& dayHex) {
    if (root.empty()) return "";
    std::vector<std::string> names;
#ifdef _WIN32
    WIN32_FIND_DATAA fd;
    HANDLE h = FindFirstFileA((root + "\\*").c_str(), &fd);
    if (h == INVALID_HANDLE_VALUE) return "";
    do { if ((fd.dwFileAttributes & FILE_ATTRIBUTE_DIRECTORY) && fd.cFileName[0] != '.') names.push_back(root + "\\" + fd.cFileName); } while (FindNextFileA(h, &fd));
    FindClose(h);
#else
    DIR* d = opendir(root.c_str());
    if (!d) return "";
    while (dirent* e = readdir(d)) if (e->d_name[0] != '.') names.push_back(root + "/" + e->d_name);
    closedir(d);
#endif
    for (const std::string& dir : names) {
        bool ok = false;
        std::string seeds = readText(dir + "/seeds.txt", ok);
        if (!ok) continue;
        char e[65] = {0}, dd[PF_HEX_CAP] = {0};
        if (pf_seeds_line(seeds.c_str(), "epoch_seed_hex", e, sizeof(e)) && pf_seeds_line(seeds.c_str(), "day_seed_hex", dd, sizeof(dd)) && epochHex == e && dayHex == dd) return dir;
    }
    return "";
}

static std::string parentDir(const std::string& p) { size_t i = p.find_last_of("/\\"); return i == std::string::npos ? "." : p.substr(0, i); }

static int runServe(Ctx& c, const Options& o, Pair* cur) {
    const uint32_t batch = 1u << o.batchLog2;
    CUdeviceptr dOut = 0;
    std::string err;
    if (c.drv.memAlloc(&dOut, (size_t)batch * 8u) != CUDA_SUCCESS) { emit("error 0 cuMemAlloc out buffer"); return 2; }
    std::vector<uint64_t> hOut(batch);
    Pair* prepared = nullptr;
    Pair* old = nullptr;
    PrepareTask* task = nullptr;
    std::string prepareRoot;   // the parent of the last prepare's pack directory: where the miner writes its packs
    emit(fmt("ready cuda %s pack %s dataset-log2 %u batch %u regs %d prepare 1 path nvrtc %d.%d driver %d.%d arch %s variant %s race %s worker %s",
             c.name.c_str(), cur->seedString.c_str(), cur->datasetLog2, batch, cur->regs, c.rtcMajor, c.rtcMinor, c.driverVersion / 1000, (c.driverVersion % 100) / 10, c.archOpt.c_str(), cur->variant.c_str(), c.race.c_str(), WORKER_VERSION));
    info(fmt("first pack %s: %s", cur->dir.c_str(), pairSummary(cur).c_str()));
    if (!cur->raceLine.empty()) emit(cur->raceLine);
    std::string line;
    while (std::getline(std::cin, line)) {
        if (line == "quit") break;
        std::vector<std::string> f = split(line);
        if (f.empty()) continue;
        // A finished prepare is reported here, between lines
        if (task && task->done) {
            task->thread.join();
            if (task->result) {
                if (prepared) releasePair(c, prepared);
                prepared = task->result;
                if (!prepared->raceLine.empty()) emit(prepared->raceLine);
                emit(fmt("prepared %s %s %.1f %s resident 2 programs 2 datasets", prepared->epochHex.c_str(), prepared->dayHex.c_str(), wallMs() - task->t0, pairSummary(prepared).c_str()));
            } else {
                emit(fmt("prepare-failed %s %s %s", task->epochHex.c_str(), task->dayHex.c_str(), task->error.c_str()));
            }
            delete task; task = nullptr;
        }
        std::string wantClass, wantEra;
        takeClassTokens(f, wantClass, wantEra);
        if (f[0] == "prepare") {
            if (f.size() < 4) { emit(fmt("prepare-failed %s %s a pack directory is needed as the third field (igneum-miner --prepare-packs <dir>)", f.size() > 1 ? f[1].c_str() : "0", f.size() > 2 ? f[2].c_str() : "0")); continue; }
            if (f[1].size() != 64) { emit(fmt("prepare-failed %s %s bad field (epoch_seed 64 hex, day_seed hex)", f[1].c_str(), f[2].c_str())); continue; }
            if (task) { emit(fmt("prepare-failed %s %s a prepare is still running", f[1].c_str(), f[2].c_str())); continue; }
            if (prepared && prepared->epochHex == f[1] && prepared->dayHex == f[2]) { emit(fmt("prepared %s %s 0 (already resident)", f[1].c_str(), f[2].c_str())); continue; }
            if (cur->epochHex == f[1] && cur->dayHex == f[2]) { emit(fmt("prepared %s %s 0 (already the current pair)", f[1].c_str(), f[2].c_str())); continue; }
            std::string dir = f[3];
            for (size_t i = 4; i < f.size(); ++i) dir += " " + f[i];   // a directory with spaces arrives as several fields
            prepareRoot = parentDir(dir);
            task = new PrepareTask();
            task->epochHex = f[1]; task->dayHex = f[2]; task->dir = dir; task->t0 = wallMs();
            task->wantClass = wantClass; task->wantEra = wantEra;
            task->thread = std::thread(prepareRun, &c, task);
            info(fmt("prepare started for epoch %.16s day %s from %s (NVRTC %s in the background)", f[1].c_str(), f[2].c_str(), dir.c_str(), c.archOpt.c_str()));
            continue;
        }
        if (f[0] != "job") { info("ignored: " + line); continue; }
        std::string jobId = f.size() > 1 ? f[1] : "0";
        if (f.size() < 8) { emit("error " + jobId + " malformed job line (need 7 fields after job)"); continue; }
        uint8_t prehash[32], epochSeed[32], daySeed[256];
        size_t pl = 0, el = 0, dl = 0;
        unsigned long long target = 0, nonceStart = 0, nonceCount = 0;
        if (!pf_unhex(f[2].c_str(), prehash, 32, &pl) || pl != 32 || std::sscanf(f[3].c_str(), "%llx", &target) != 1 ||
            std::sscanf(f[4].c_str(), "%llu", &nonceStart) != 1 || std::sscanf(f[5].c_str(), "%llu", &nonceCount) != 1 ||
            !pf_unhex(f[6].c_str(), epochSeed, 32, &el) || el != 32 || !pf_unhex(f[7].c_str(), daySeed, sizeof(daySeed), &dl)) {
            emit("error " + jobId + " bad field (prehash 64 hex, target 16 hex, nonce_start u64, nonce_count u64, epoch_seed 64 hex, day_seed hex)"); continue;
        }
        if (nonceCount == 0 || nonceCount % 32 != 0 || (nonceStart & 31) != 0) { emit("error " + jobId + " nonce_start must be 32-aligned and nonce_count a non-zero multiple of 32"); continue; }
        uint32_t sw[8], kw[8];
        pf_seed_words_from_bytes(epochSeed, 32, sw);
        pf_seed_words_from_bytes(daySeed, dl, kw);
        double t0 = wallMs();
        bool switched = false;
        if (!pairIsClass(cur, f[6], f[7], wantClass, wantEra) && !task && !pairIsClass(prepared, f[6], f[7], wantClass, wantEra)) {
            // Self-heal: a job on seeds this worker has no pair for and no prepare in flight (a prepare failed, or
            // the miner never sent one). The miner writes a pack per pair under its --prepare-packs root; find it by
            // seeds.txt and build it now, in the foreground. The miner only re-sends prepare for the pair after this one.
            std::string dir = findPackFor(prepareRoot, f[6], f[7]);
            if (dir.empty()) dir = findPackFor(parentDir(o.pack) + "/prepare", f[6], f[7]);
            if (dir.empty()) dir = findPackFor(parentDir(o.pack), f[6], f[7]);
            if (!dir.empty()) {
                info(fmt("job %s is for epoch %.16s day %s, which is not resident; building its pack %s now (foreground)", jobId.c_str(), f[6].c_str(), f[7].c_str(), dir.c_str()));
                std::string berr;
                Pair* p = buildPair(c, dir, nullptr, berr, true);
                if (p && (p->epochHex != f[6] || p->dayHex != f[7])) { berr = "the pack in " + dir + " is for other seeds"; releasePair(c, p); p = nullptr; }
                if (p) { if (prepared) releasePair(c, prepared); prepared = p; info(fmt("built %s: %s", dir.c_str(), pairSummary(p).c_str())); if (!p->raceLine.empty()) emit(p->raceLine); }
                else emit("error " + jobId + " could not build " + dir + ": " + berr);
            }
        }
        if (!pairIsClass(cur, f[6], f[7], wantClass, wantEra)) {
            char why[256];
            if (pairIsClass(prepared, f[6], f[7], wantClass, wantEra)) {
                if (old) releasePair(c, old);
                old = cur; cur = prepared; prepared = nullptr; switched = true;
                info(fmt("switched to the prepared pair epoch %.16s day %s (class %s) in %.2f ms", cur->epochHex.c_str(), cur->dayHex.c_str(), cur->programClass.c_str(), wallMs() - t0));
            } else if (pairIs(cur, f[6], f[7]) && !pf_pack_class_ok(cur->programClass.c_str(), cur->eraHex.c_str(), wantClass.c_str(), wantEra.c_str(), why, sizeof(why))) {
                // Counter ASIC 2.0: right seeds, wrong class or era. The miner prepares the pair again from a pack of
                // the class the chain is on (the `need` line), and the pack of the other class is never mined.
                emit(fmt("need %s %s", f[6].c_str(), f[7].c_str()));
                emit(fmt("error %s pack %s: %s", jobId.c_str(), cur->dir.c_str(), why));
                continue;
            } else if (!hexEq(cur->epochHex, f[6])) {
                emit(fmt("need %s %s", f[6].c_str(), f[7].c_str()));   // the miner prepares this pair (4 October 2026)
                emit(fmt("error %s epoch seed mismatch: this worker holds epoch %.16s (program words %08x %08x ...)%s, the job is for epoch %.16s (bare seed words %08x %08x ...); send prepare with a pack directory",
                         jobId.c_str(), cur->epochHex.c_str(), cur->sw[0], cur->sw[1], prepared ? " plus one prepared pair" : "", f[6].c_str(), sw[0], sw[1]));
                continue;
            } else {
                emit(fmt("need %s %s", f[6].c_str(), f[7].c_str()));
                emit(fmt("error %s day seed mismatch: this worker's cache is for key %08x %08x ..., the job's day seed %s gives %08x %08x ...; send prepare with a pack directory",
                         jobId.c_str(), cur->kw[0], cur->kw[1], f[7].c_str(), kw[0], kw[1]));
                continue;
            }
        }
        uint64_t remaining = nonceCount, hashes = 0;
        uint32_t hi = (uint32_t)(nonceStart >> 32), lo = (uint32_t)nonceStart;
        bool failed = false;
        while (remaining > 0) {
            uint64_t room = (uint64_t)(0xffffffffu - lo) + 1ull;
            uint64_t chunk64 = remaining < batch ? remaining : batch;
            if (chunk64 > room) chunk64 = room;
            uint32_t chunk = (uint32_t)chunk64;
            uint32_t iw[8];
            {
                uint8_t b[49];
                std::memcpy(b, "igneum-block/", 13);
                std::memcpy(b + 13, prehash, 32);
                b[45] = (uint8_t)hi; b[46] = (uint8_t)(hi >> 8); b[47] = (uint8_t)(hi >> 16); b[48] = (uint8_t)(hi >> 24);
                pf_seed_words_from_bytes(b, 49, iw);
            }
            // A chunk that is not a multiple of the block is finished one 32-lane block at a time. The block is the
            // pair's (its winning variant's). The mutex gives a race its exclusive windows between chunks.
            std::lock_guard<std::mutex> hold(gpuMutex);
            uint32_t block = 32u * (uint32_t)cur->blockWarps;
            uint32_t main = chunk - (chunk % block);
            CUresult r = CUDA_SUCCESS;
            if (main > 0 && !launchHash(c, cur, dOut, lo, iw, main, block, nullptr, err)) { emit("error " + jobId + " dispatch failed: " + err); failed = true; break; }
            if (main < chunk) {
                for (uint32_t off = main; off < chunk && !failed; off += 32u) if (!launchHash(c, cur, dOut + (CUdeviceptr)off * 8u, lo + off, iw, 32u, 32u, nullptr, err)) { emit("error " + jobId + " dispatch failed: " + err); failed = true; }
                if (failed) break;
            }
            r = c.drv.streamSynchronize(nullptr);
            if (r == CUDA_SUCCESS) r = c.drv.memcpyDtoH(hOut.data(), dOut, (size_t)chunk * 8u);
            if (r != CUDA_SUCCESS) { emit("error " + jobId + " dispatch failed: " + c.err(r)); failed = true; break; }
            for (uint32_t i = 0; i < chunk; ++i) if (hOut[i] <= target) {
                uint64_t nonce = ((uint64_t)hi << 32) | (uint64_t)(uint32_t)(lo + i);
                std::printf("found %s %llu %016llx\n", jobId.c_str(), (unsigned long long)nonce, (unsigned long long)hOut[i]);
            }
            std::fflush(stdout);
            hashes += chunk;
            remaining -= chunk;
            if (chunk64 == room) { hi += 1u; lo = 0u; } else lo += chunk;
        }
        if (failed) continue;
        emit(fmt("done %s %llu %.2f", jobId.c_str(), (unsigned long long)hashes, wallMs() - t0));
        if (switched && old) { releasePair(c, old); old = nullptr; info("dropped the previous pair (its program, cache and dataset)"); }
    }
    if (task) { task->thread.join(); if (task->result) releasePair(c, task->result); delete task; }
    c.drv.memFree(dOut);
    if (old) releasePair(c, old);
    if (prepared) releasePair(c, prepared);
    releasePair(c, cur);
    return 0;
}

// ---------------------------------------------------------------------------------------------
// Main

// ---------------------------------------------------------------------------------------------
// Read-width experiment (5 October 2026, docs/plans/read-width.md): --bench and --memprobe

static uint64_t fnv1a64Bytes(const void* p, size_t n) {
    const uint8_t* b = (const uint8_t*)p;
    uint64_t h = 0xcbf29ce484222325ull;
    for (size_t i = 0; i < n; ++i) { h ^= b[i]; h *= 0x100000001b3ull; }
    return h;
}

// --bench: the pair is built and self-tested (vectors through the bound kernel with the seed words); then a warm-up
// dispatch at base nonce 0 (fingerprinted) and --batches timed dispatches of 2^B nonces, wall time around
// cuStreamSynchronize (the driver API path loads no event symbols; a 2^24 dispatch is 60 to 900 ms on the cards here,
// so the launch overhead is under 1 percent).
static int runBench(Ctx& c, const Options& o, Pair* p) {
    uint32_t nonces = 1u << o.batchLog2, block = 32u * (uint32_t)c.blockWarps;
    if (p->persistent) { uint32_t unit = 32u * (uint32_t)p->warps; nonces = (nonces / unit) * unit; if (nonces == 0) nonces = unit; }
    CUdeviceptr dOut = 0;
    std::string err;
    if (c.drv.memAlloc(&dOut, (size_t)nonces * 8u) != CUDA_SUCCESS) { std::printf("FAIL: cuMemAlloc out\n"); return 2; }
    std::vector<uint64_t> hOut(nonces);
    double sum = 0, warm = 0;
    uint64_t fp = 0;
    for (int b = -1; b < o.batches; ++b) {
        double t0 = wallMs();
        if (!launchHash(c, p, dOut, (uint32_t)(b + 1) * nonces, p->sw, nonces, block, nullptr, err)) { std::printf("FAIL: %s\n", err.c_str()); return 2; }
        CUresult r = c.drv.streamSynchronize(nullptr);
        if (r != CUDA_SUCCESS) { std::printf("FAIL: dispatch %d: %s\n", b, c.err(r).c_str()); return 2; }
        double ms = wallMs() - t0;
        if (b < 0) {
            warm = ms;
            if (c.drv.memcpyDtoH(hOut.data(), dOut, (size_t)nonces * 8u) != CUDA_SUCCESS) { std::printf("FAIL: read-back\n"); return 2; }
            fp = fnv1a64Bytes(hOut.data(), (size_t)nonces * 8u);
        } else sum += ms;
    }
    c.drv.memFree(dOut);
    std::string dev = c.name; for (char& ch : dev) if (ch == ' ') ch = '_';
    std::printf("warm-up dispatch (base 0): %.2f ms; %d timed dispatches of %u nonces: mean %.2f ms\n", warm, o.batches, nonces, sum / o.batches);
    std::printf("RESULT pack=%s class=%s device=%s arch=%s regs=%d blocks_per_sm=%d warps=%d resident=%d arena_mib=%llu hot_mib=%u hot_slots=%u hot_fill_ms=%.2f nonces=%u batches=%d check=%s fingerprint=%016llx mhs=%.3f loads=%u bytes=%u scratch_ops=%u time=wall\n",
                p->dir.c_str(), p->loadClass.c_str(), dev.c_str(), c.archOpt.c_str(), p->regs, p->blocksPerSM, p->warps, p->residentWarps, (unsigned long long)(p->scratchBytes >> 20), p->hotMb, p->hotSlots, p->hotMs, nonces, o.batches,
                p->checked ? (p->checkPass ? "PASS" : "FAIL") : "skipped", (unsigned long long)fp, (double)nonces * (double)o.batches / (sum / 1000.0) / 1e6,
                p->loadsPerHash, p->bytesPerHash, p->scratchOps * 8u);
    return 0;
}

// --memprobe: the OpenCL worker's table (proto-opencl/host.c, 5 October 2026) in CUDA C through NVRTC, so the two
// vendors are probed with the same access patterns: a dependent chain of random 4-byte reads (the hash's pattern),
// eight independent chains per lane, dependent random 16-byte and 64-byte reads, a coalesced stream and an integer
// chain, at 4, 64 and 1024 MiB. Wall time around cuStreamSynchronize, best of 3, a fresh seed per repetition.
static const char* PROBE_CUDA =
    "#include <cstdint>\n"
    "__device__ __forceinline__ uint32_t pm_mix(uint32_t x) { x ^= x >> 16; x *= 0x7feb352du; x ^= x >> 15; x *= 0x846ca68bu; x ^= x >> 16; return x; }\n"
    "extern \"C\" __global__ void probe_fill(uint32_t* ds, uint32_t n) { uint32_t i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) ds[i] = pm_mix(i ^ 0x9E3779B9u); }\n"
    "extern \"C\" __global__ void probe_chase(const uint32_t* ds, uint32_t mask, uint32_t steps, uint32_t seed, uint32_t* out) {\n"
    "    uint32_t g = blockIdx.x * blockDim.x + threadIdx.x; uint32_t x = pm_mix(g ^ seed);\n"
    "    for (uint32_t s = 0u; s < steps; ++s) x = ds[x & mask] ^ (x * 0x9E3779B1u + s);\n"
    "    out[g] = x;\n"
    "}\n"
    "extern \"C\" __global__ void probe_indep(const uint32_t* ds, uint32_t mask, uint32_t steps, uint32_t seed, uint32_t* out) {\n"
    "    uint32_t g = blockIdx.x * blockDim.x + threadIdx.x;\n"
    "    uint32_t x0 = pm_mix(g * 8u ^ seed), x1 = pm_mix((g * 8u + 1u) ^ seed), x2 = pm_mix((g * 8u + 2u) ^ seed), x3 = pm_mix((g * 8u + 3u) ^ seed);\n"
    "    uint32_t x4 = pm_mix((g * 8u + 4u) ^ seed), x5 = pm_mix((g * 8u + 5u) ^ seed), x6 = pm_mix((g * 8u + 6u) ^ seed), x7 = pm_mix((g * 8u + 7u) ^ seed);\n"
    "    for (uint32_t s = 0u; s < steps; ++s) {\n"
    "        x0 = ds[x0 & mask] ^ (x0 * 0x9E3779B1u + s); x1 = ds[x1 & mask] ^ (x1 * 0x9E3779B1u + s);\n"
    "        x2 = ds[x2 & mask] ^ (x2 * 0x9E3779B1u + s); x3 = ds[x3 & mask] ^ (x3 * 0x9E3779B1u + s);\n"
    "        x4 = ds[x4 & mask] ^ (x4 * 0x9E3779B1u + s); x5 = ds[x5 & mask] ^ (x5 * 0x9E3779B1u + s);\n"
    "        x6 = ds[x6 & mask] ^ (x6 * 0x9E3779B1u + s); x7 = ds[x7 & mask] ^ (x7 * 0x9E3779B1u + s);\n"
    "    }\n"
    "    out[g] = x0 ^ x1 ^ x2 ^ x3 ^ x4 ^ x5 ^ x6 ^ x7;\n"
    "}\n"
    "extern \"C\" __global__ void probe_line16(const uint4* ds, uint32_t vecMask, uint32_t steps, uint32_t seed, uint32_t* out) {\n"
    "    uint32_t g = blockIdx.x * blockDim.x + threadIdx.x; uint32_t x = pm_mix(g ^ seed);\n"
    "    for (uint32_t s = 0u; s < steps; ++s) { uint4 a = ds[x & vecMask]; x = (a.x ^ a.y ^ a.z ^ a.w) ^ (x * 0x9E3779B1u + s); }\n"
    "    out[g] = x;\n"
    "}\n"
    "extern \"C\" __global__ void probe_line(const uint4* ds, uint32_t lineMask, uint32_t steps, uint32_t seed, uint32_t* out) {\n"
    "    uint32_t g = blockIdx.x * blockDim.x + threadIdx.x; uint32_t x = pm_mix(g ^ seed);\n"
    "    for (uint32_t s = 0u; s < steps; ++s) { uint32_t l = (x & lineMask) * 4u; uint4 a = ds[l], b = ds[l + 1u], c = ds[l + 2u], d = ds[l + 3u]; x = (a.x ^ b.y ^ c.z ^ d.w) ^ (x * 0x9E3779B1u + s); }\n"
    "    out[g] = x;\n"
    "}\n"
    "extern \"C\" __global__ void probe_stream(const uint4* ds, uint32_t perLane, uint32_t* out) {\n"
    "    uint32_t g = blockIdx.x * blockDim.x + threadIdx.x, n = gridDim.x * blockDim.x; uint4 acc = make_uint4(0u, 0u, 0u, 0u);\n"
    "    for (uint32_t s = 0u; s < perLane; ++s) { uint4 v = ds[s * n + g]; acc.x ^= v.x; acc.y ^= v.y; acc.z ^= v.z; acc.w ^= v.w; }\n"
    "    out[g] = acc.x ^ acc.y ^ acc.z ^ acc.w;\n"
    "}\n"
    "extern \"C\" __global__ void probe_alu(uint32_t steps, uint32_t seed, uint32_t* out) {\n"
    "    uint32_t g = blockIdx.x * blockDim.x + threadIdx.x; uint32_t x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u;\n"
    "    for (uint32_t s = 0u; s < steps; ++s) { x = x * 0x9E3779B1u + ((y << 7u) | (y >> 25u)); y = (y ^ x) + s; }\n"
    "    out[g] = x ^ y;\n"
    "}\n";

static double probeLaunch(Ctx& c, CUfunction f, size_t lanes, size_t local, int reps, int seedArg, uint32_t seed, void** args) {
    double best = -1;
    for (int r = 0; r < reps; ++r) {
        uint32_t s = seed + (uint32_t)r * 0x9E3779B9u;
        if (seedArg >= 0) args[seedArg] = &s;
        double t0 = wallMs();
        if (c.drv.launchKernel(f, (unsigned)(lanes / local), 1, 1, (unsigned)local, 1, 1, 0, nullptr, args, nullptr) != CUDA_SUCCESS) return -1;
        if (c.drv.streamSynchronize(nullptr) != CUDA_SUCCESS) return -1;
        double ms = wallMs() - t0;
        if (best < 0 || ms < best) best = ms;
    }
    return best;
}

static int runMemprobe(Ctx& c, const Options& o) {
    Compiled cp;
    std::string err;
    if (!rtcCompile(c, PROBE_CUDA, "probe.cu", "", "", {}, cp, err)) { std::printf("memprobe: build FAILED: %s\n", err.c_str()); return 2; }
    CUmodule mod = nullptr;
    if (c.drv.moduleLoadData(&mod, cp.image.data()) != CUDA_SUCCESS) { std::printf("memprobe: cuModuleLoadData failed\n"); return 2; }
    CUfunction kFill, kChase, kIndep, kLine16, kLine, kStream, kAlu;
    const char* names[7] = { "probe_fill", "probe_chase", "probe_indep", "probe_line16", "probe_line", "probe_stream", "probe_alu" };
    CUfunction* fns[7] = { &kFill, &kChase, &kIndep, &kLine16, &kLine, &kStream, &kAlu };
    for (int i = 0; i < 7; ++i) if (c.drv.moduleGetFunction(fns[i], mod, names[i]) != CUDA_SUCCESS) { std::printf("memprobe: %s not in the module\n", names[i]); return 2; }
    int sizes[3] = { 4, 64, 1024 }, nSizes = 3;
    if (o.probeMib > 0) { sizes[0] = o.probeMib; nSizes = 1; }
    const size_t lanesList[8] = { 256, 1024, 1u << 12, 1u << 14, 1u << 16, 1u << 18, 1u << 20, 1u << 22 };
    const size_t groups[2] = { 32, 256 };
    const uint32_t STEPS = 256u, ALU_STEPS = 4096u;
    const size_t maxLanes = 1u << 22;
    CUdeviceptr dOut = 0;
    if (c.drv.memAlloc(&dOut, maxLanes * 4u) != CUDA_SUCCESS) { std::printf("memprobe: cuMemAlloc out\n"); return 2; }
    std::printf("memprobe on %s (sm_%d%d, %d SMs, driver %d.%d, NVRTC %d.%d), wall time around cuStreamSynchronize\n", c.name.c_str(), c.major, c.minor, c.sms, c.driverVersion / 1000, (c.driverVersion % 100) / 10, c.rtcMajor, c.rtcMinor);
    std::printf("| probe | MiB | block | lanes in flight | steps per lane | best ms | G loads/s | ns per dependent load |\n|---|---|---|---|---|---|---|---|\n");
    for (int si = 0; si < nSizes; ++si) {
        int mib = sizes[si];
        uint64_t bytes = (uint64_t)mib << 20;
        uint32_t words = (uint32_t)(bytes / 4ull), mask = words - 1u, n = words;
        CUdeviceptr dDs = 0;
        if (c.drv.memAlloc(&dDs, (size_t)bytes) != CUDA_SUCCESS) { std::printf("| chase | %d | skipped: cuMemAlloc failed | | | | | |\n", mib); continue; }
        { void* a[2] = { &dDs, &n }; probeLaunch(c, kFill, ((size_t)words + 255) / 256 * 256, 256, 1, -1, 0, a); }
        for (int gi = 0; gi < 2; ++gi) {
            size_t local = groups[gi];
            for (int li = 0; li < 8; ++li) {
                size_t lanes = lanesList[li];
                if (lanes < local) continue;
                uint32_t seed = 0x1234567u + (uint32_t)li * 977u, steps = STEPS;
                void* a[5] = { &dDs, &mask, &steps, &seed, &dOut };
                double ms = probeLaunch(c, kChase, lanes, local, 3, 3, seed, a);
                std::printf("| chase | %d | %zu | %zu | %u | %.3f | %.3f | %.0f |\n", mib, local, lanes, STEPS, ms, (double)lanes * STEPS / (ms / 1000.0) / 1e9, ms * 1e6 / STEPS);
                std::fflush(stdout);
            }
        }
        for (size_t lanes = 1u << 16; lanes <= maxLanes; lanes <<= 2) {
            uint32_t seed = 0x7654321u, steps = STEPS;
            void* a[5] = { &dDs, &mask, &steps, &seed, &dOut };
            double ms = probeLaunch(c, kIndep, lanes, 256, 3, 3, seed, a);
            std::printf("| indep x8 | %d | 256 | %zu | %u | %.3f | %.3f | (8 loads in flight per lane) |\n", mib, lanes, STEPS, ms, (double)lanes * 8.0 * STEPS / (ms / 1000.0) / 1e9);
        }
        for (size_t lanes = 1u << 14; lanes <= maxLanes; lanes <<= 2) {
            uint32_t vecMask = (words / 4u) - 1u, seed = 0x2718281u, steps = STEPS;
            void* a[5] = { &dDs, &vecMask, &steps, &seed, &dOut };
            double ms = probeLaunch(c, kLine16, lanes, 256, 3, 3, seed, a);
            std::printf("| line 16 B | %d | 256 | %zu | %u | %.3f | %.3f G reads/s | %.1f GB/s in 16 B reads |\n", mib, lanes, STEPS, ms, (double)lanes * STEPS / (ms / 1000.0) / 1e9, (double)lanes * STEPS * 16.0 / (ms / 1000.0) / 1e9);
        }
        for (size_t lanes = 1u << 14; lanes <= maxLanes; lanes <<= 2) {
            uint32_t lineMask = (words / 16u) - 1u, seed = 0x3141592u, steps = STEPS;
            void* a[5] = { &dDs, &lineMask, &steps, &seed, &dOut };
            double ms = probeLaunch(c, kLine, lanes, 256, 3, 3, seed, a);
            std::printf("| line 64 B | %d | 256 | %zu | %u | %.3f | %.3f G lines/s | %.1f GB/s in lines |\n", mib, lanes, STEPS, ms, (double)lanes * STEPS / (ms / 1000.0) / 1e9, (double)lanes * STEPS * 64.0 / (ms / 1000.0) / 1e9);
        }
        {
            size_t lanes = 1u << 20;
            uint32_t perLane = (uint32_t)((uint64_t)words / 4ull / (uint64_t)lanes);
            if (perLane == 0) { perLane = 1; lanes = (size_t)words / 4u; }
            double bytesRead = (double)perLane * (double)lanes * 16.0;
            void* a[3] = { &dDs, &perLane, &dOut };
            double ms = probeLaunch(c, kStream, lanes, 256, 3, -1, 0, a);
            std::printf("| stream | %d | 256 | %zu | %u | %.3f | %.1f GB/s coalesced | (%.0f MiB read once) |\n", mib, lanes, perLane, ms, bytesRead / (ms / 1000.0) / 1e9, bytesRead / 1048576.0);
        }
        c.drv.memFree(dDs);
        std::fflush(stdout);
    }
    {
        size_t lanes = 1u << 20;
        uint32_t seed = 0x2468aceu, steps = ALU_STEPS;
        void* a[3] = { &steps, &seed, &dOut };
        double ms = probeLaunch(c, kAlu, lanes, 256, 3, 1, seed, a);
        double ops = (double)lanes * ALU_STEPS * 5.0;
        std::printf("| alu | 0 | 256 | %zu | %u | %.3f | %.1f G int ops/s | %.3f G steps/s per SM (approximate: 5 ops per step counted) |\n", lanes, ALU_STEPS, ms, ops / (ms / 1000.0) / 1e9, (double)lanes * ALU_STEPS / (ms / 1000.0) / 1e9 / (c.sms ? c.sms : 1));
    }
    c.drv.memFree(dOut);
    c.drv.moduleUnload(mod);
    std::printf("memprobe: done\n");
    return 0;
}

int main(int argc, char** argv) {
    Options o = parseArgs(argc, argv);
    Ctx c;
    c.blockWarps = o.blockWarps;
    c.race = o.race; c.raceBenchMs = o.raceBenchMs; c.raceBudgetS = o.raceBudgetS; c.raceRounds = o.raceRounds; c.batchLog2 = o.batchLog2; c.pinned = o.pinned;
    c.warps = o.warps; c.batches = o.batches;
    if (o.bench || o.memprobe) c.race = "off";
    if (!o.tuningPath.empty()) { bool ok = false; c.tuning = readText(o.tuningPath, ok); if (!ok) c.tuning.clear(); }
    std::string err, drvLib, rtcLib;
    if (!loadDriver(c.drv, err, drvLib)) { emit("error 0 " + err); return 2; }
    if (!loadNvrtc(c.rtc, err, rtcLib)) { emit("error 0 " + err); return 2; }
    if (!openDevice(c, o.device, o.arch, err)) { emit("error 0 " + err); return 2; }
    info(fmt("igneum-worker-cuda %s: device %d %s (sm_%d%d, %d SMs), driver %d.%d from %s, NVRTC %d.%d from %s, target %s (%s)",
             WORKER_VERSION, o.device, c.name.c_str(), c.major, c.minor, c.sms, c.driverVersion / 1000, (c.driverVersion % 100) / 10, drvLib.c_str(), c.rtcMajor, c.rtcMinor, rtcLib.c_str(), c.archOpt.c_str(), c.why.c_str()));
    if (!c.tuning.empty()) info(fmt("tuning file %s (%zu bytes): %s", o.tuningPath.c_str(), c.tuning.size(), readTuning(c.tuning, c.name).found ? "has an entry for this card" : "no entry for this card"));
    if (o.memprobe) { int rc = runMemprobe(c, o); c.drv.primaryCtxRelease(c.dev); return rc; }
    double t0 = wallMs();
    Pair* cur = buildPair(c, o.pack, nullptr, err, !o.check && !o.bench);
    if (!cur) { emit("error 0 " + err); return 1; }
    if (o.bench) {
        std::printf("pack %s on %s: %s\n", o.pack.c_str(), c.name.c_str(), pairSummary(cur).c_str());
        int rc = runBench(c, o, cur);
        releasePair(c, cur);
        c.drv.primaryCtxRelease(c.dev);
        return rc;
    }
    if (o.raceOnly) {
        std::printf("race %s on %s (%s, %d SMs, driver %d.%d, NVRTC %d.%d, %s): %s\n", o.pack.c_str(), c.name.c_str(), c.archOpt.c_str(), c.sms, c.driverVersion / 1000, (c.driverVersion % 100) / 10, c.rtcMajor, c.rtcMinor, c.why.c_str(), pairSummary(cur).c_str());
        std::printf("%s\n", cur->raceLine.c_str());
        std::printf("winner %s: %d registers, %d blocks/SM at %d warp(s)/block\n", cur->variant.c_str(), cur->regs, cur->blocksPerSM, cur->blockWarps);
        releasePair(c, cur);
        return 0;
    }
    if (o.check) {
        std::printf("check PASS %s in %.0f ms: %s\n", o.pack.c_str(), wallMs() - t0, pairSummary(cur).c_str());
        std::printf("  epoch %s day %s, dataset %s words, cache 2^%u words in %u segments, %d registers, %d blocks/SM at %d warp(s)/block, target %s\n",
                    cur->epochHex.c_str(), cur->dayHex.c_str(), (cur->mulshift ? fmt("%u (%u items, multiply-shift)", cur->words, cur->items) : fmt("2^%u", cur->datasetLog2)).c_str(), (unsigned)__builtin_ctz(cur->cacheWords), cur->cacheSegments, cur->regs, cur->blocksPerSM, c.blockWarps, c.archOpt.c_str());
        releasePair(c, cur);
        return 0;
    }
    int rc = runServe(c, o, cur);
    c.drv.primaryCtxRelease(c.dev);
    return rc;
}