// 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.
//
// Usage: igneum-worker-cuda --serve --pack <dir> [--device D] [--batch-log2 22] [--block-warps 1] [--arch sm_120|auto]
//        igneum-worker-cuda --check --pack <dir> [--device D]      compile, build, self-test, print timings, exit 0/1

#include <cstdint>
#include <cstdarg>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <chrono>
#include <string>
#include <vector>
#include <thread>
#include <atomic>
#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, 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;
    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;
    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) {
    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.
    const char* opts[3] = { archOpt.c_str(), "--std=c++17", "-default-device" };
    r = c.rtc.compileProgram(prog, 3, opts);
    {
        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

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;
    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;
};

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->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);
    void* args[5] = { &p->ds, &out, &baseNonce, &mask, &a };
    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;
}

// Compiles the pack in `dir`, builds its cache and dataset on stream `s`, runs the self-test. 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) {
    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);
    p->datasetLog2 = pk.datasetLog2; p->words = 1u << pk.datasetLog2; p->cacheWords = 1u << pk.cacheLog2Words; p->cacheSegments = pk.cacheSegments;
    // Compile
    Compiled ck, cb;
    if (!rtcCompile(c, kernelDev, "kernel.cu", programH, memhardH, { "igneum_cache_fill", "igneum_build" }, 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; }
        c.drv.funcGetAttribute(&p->regs, CU_FUNC_ATTRIBUTE_NUM_REGS, p->fHashBound);
        c.drv.occupancy(&p->blocksPerSM, p->fHashBound, 32 * c.blockWarps, 0);
    }
    // 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 + (64u << 20)) {
            err = fmt("%llu MiB free on the device, this pack needs %llu MiB (cache %llu + dataset %llu)", (unsigned long long)(freeB >> 20), (unsigned long long)((cacheBytes + dsBytes) >> 20), (unsigned long long)(cacheBytes >> 20), (unsigned long long)(dsBytes >> 20));
            releasePair(c, p); return nullptr;
        }
    }
    {
        CUresult r = c.drv.memAlloc(&p->cache, cacheBytes);
        if (r != CUDA_SUCCESS) { err = "cuMemAlloc cache: " + c.err(r); p->cache = 0; 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); 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; releasePair(c, p); return nullptr; }
        uint32_t nItems = p->words / 16u, block = 256u, grid = (nItems + block - 1u) / block;
        void* args[3] = { &p->ds, &p->cache, &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);
        if (r != CUDA_SUCCESS) { err = "dataset build: " + c.err(r); releasePair(c, p); return nullptr; }
    }
    p->dsMs = 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);
        char line[1024];
        p->checkPass = pf_selftest(&pk, cacheHead, cacheLast, fnv, dsHead, dsLast, samples.data(), vec.data(), 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; }
    return p;
}

static std::string pairSummary(const Pair* p) {
    return fmt("nvrtc %.0f cache %.0f dataset %.0f check %.0f ms; %s", p->compileMs, p->cacheMs, p->dsMs, p->checkMs, p->check.c_str());
}

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

struct PrepareTask {
    std::string epochHex, dayHex, dir, error;
    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);
    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;
    }
    t->result = p;
    t->error = err;
    t->done = true;
}

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

struct Options {
    bool serve = false, check = false;
    int device = 0, batchLog2 = 22, blockWarps = 1;
    std::string pack, arch = "auto";
};

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", 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 == "--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) { usage(); std::exit(2); }
    if (o.pack.empty()) { 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 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(), WORKER_VERSION));
    info(fmt("first pack %s: %s", cur->dir.c_str(), pairSummary(cur).c_str()));
    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;
                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;
        }
        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->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 ((std::memcmp(sw, cur->sw, 32) != 0 || std::memcmp(kw, cur->kw, 32) != 0) && !task &&
            !(prepared && std::memcmp(sw, prepared->sw, 32) == 0 && std::memcmp(kw, prepared->kw, 32) == 0)) {
            // 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);
                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())); }
                else emit("error " + jobId + " could not build " + dir + ": " + berr);
            }
        }
        if (std::memcmp(sw, cur->sw, 32) != 0 || std::memcmp(kw, cur->kw, 32) != 0) {
            if (prepared && std::memcmp(sw, prepared->sw, 32) == 0 && std::memcmp(kw, prepared->kw, 32) == 0) {
                if (old) releasePair(c, old);
                old = cur; cur = prepared; prepared = nullptr; switched = true;
                info(fmt("switched to the prepared pair epoch %.16s day %s in %.2f ms", cur->epochHex.c_str(), cur->dayHex.c_str(), wallMs() - t0));
            } else if (std::memcmp(sw, cur->sw, 32) != 0) {
                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 (seed words %08x %08x ...)%s, the job's epoch seed %.16s gives %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
            uint32_t block = 32u * (uint32_t)o.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

int main(int argc, char** argv) {
    Options o = parseArgs(argc, argv);
    Ctx c;
    c.blockWarps = o.blockWarps;
    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()));
    double t0 = wallMs();
    Pair* cur = buildPair(c, o.pack, nullptr, err);
    if (!cur) { emit("error 0 " + err); return 1; }
    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 2^%u 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->datasetLog2, (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;
}