// igneum-bench-cl: OpenCL host program for Igneum's random-program proof-of-work test harness.
// The portable third path after Apple Metal (proto-metal) and NVIDIA CUDA (proto-cuda): it runs on AMD (Windows
// and Linux), NVIDIA, Intel and, as a correctness check only, on Apple's deprecated OpenCL 1.2 runtime.
//
// TEST HARNESS ONLY. No pool, no network, no wallet, no mining protocol. It fills the dataset on the device,
// checks the device against vectors produced on the Mac (proto-metal), and times the kernel.
//
// C99 plus the OpenCL 1.2 API, nothing else. The kernels are compiled from packs/<seed>/kernel.cl at runtime.
// The pack's program.h, vectors.h and (memory-hard packs) memhard.h are included at compile time; memhard.h is the
// host reference that fills the cache on one thread and derives dataset words for the self-test.
//
// Build: see README.md (macOS -framework OpenCL, Linux -lOpenCL, Windows cl.exe + OpenCL.lib), or build.sh / build.bat.

#define _CRT_SECURE_NO_WARNINGS
#define CL_TARGET_OPENCL_VERSION 120
#define CL_USE_DEPRECATED_OPENCL_1_2_APIS
#if defined(__APPLE__) && !defined(IGNEUM_KHR_HEADERS)
#include <OpenCL/opencl.h>
#else
#include <CL/cl.h>
#endif
#ifdef IGNEUM_CL_DYNAMIC
#include "cl_dynamic.h"   /* Windows one-click build: OpenCL.dll loaded at run time, no import library */
#endif

#include <stdint.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#ifdef _WIN32
#define WIN32_LEAN_AND_MEAN
#include <windows.h>
#define strtok_r strtok_s
#else
#include <time.h>
#include <dlfcn.h>
#include <pthread.h>
#endif

#define IGNEUM_NO_CUDA
#include "program.h"
#include "vectors.h"
#include "../proto-cuda/nvrtc/packfile.h"   /* --pack: a pack read at run time (generic serve mode, 4 October 2026) */

#ifndef IGNEUM_DATASET_MODE
#define IGNEUM_DATASET_MODE 0
#endif
#if IGNEUM_DATASET_MODE == 1
#include "memhard.h"
#endif
#ifndef IGNEUM_KERNEL_PATH
#define IGNEUM_KERNEL_PATH "kernel.cl"
#endif

// Sub-group query constants (cl_khr_subgroups / OpenCL 2.1). Spelled out because OpenCL 1.2 headers lack them.
#define IG_CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE 0x2033
#define IG_CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE 0x2034
// Vendor device attributes (cl_amd_device_attribute_query, cl_nv_device_attribute_query).
#define IG_CL_DEVICE_WAVEFRONT_WIDTH_AMD 0x4043
#define IG_CL_DEVICE_WARP_SIZE_NV 0x4003
// Where the card sits on the PCI bus, so the app can tell one physical card listed by two OpenCL platforms (two
// AMD ICDs after a driver upgrade, PC 1 on 5 October 2026) from two cards: CL_DEVICE_TOPOLOGY_AMD and the NVIDIA pair.
#define IG_CL_DEVICE_TOPOLOGY_AMD 0x4037
#define IG_CL_DEVICE_TOPOLOGY_TYPE_PCIE_AMD 1
#define IG_CL_DEVICE_PCI_BUS_ID_NV 0x4008
#define IG_CL_DEVICE_PCI_SLOT_ID_NV 0x4009
typedef union {
    struct { cl_uint type; cl_uint data[5]; } raw;
    struct { cl_uint type; cl_char unused[17]; cl_char bus; cl_char device; cl_char function; } pcie;
} ig_topology_amd;

typedef cl_int (CL_API_CALL *ig_pfn_subgroup_info)(cl_kernel, cl_device_id, cl_uint, size_t, const void*, size_t, void*, size_t*);

// ---------------------------------------------------------------------------------------------
// Errors and timing

static const char* clErrName(cl_int e) {
    switch (e) {
    case CL_SUCCESS: return "CL_SUCCESS";
    case CL_DEVICE_NOT_FOUND: return "CL_DEVICE_NOT_FOUND";
    case CL_DEVICE_NOT_AVAILABLE: return "CL_DEVICE_NOT_AVAILABLE";
    case CL_COMPILER_NOT_AVAILABLE: return "CL_COMPILER_NOT_AVAILABLE";
    case CL_MEM_OBJECT_ALLOCATION_FAILURE: return "CL_MEM_OBJECT_ALLOCATION_FAILURE";
    case CL_OUT_OF_RESOURCES: return "CL_OUT_OF_RESOURCES";
    case CL_OUT_OF_HOST_MEMORY: return "CL_OUT_OF_HOST_MEMORY";
    case CL_PROFILING_INFO_NOT_AVAILABLE: return "CL_PROFILING_INFO_NOT_AVAILABLE";
    case CL_BUILD_PROGRAM_FAILURE: return "CL_BUILD_PROGRAM_FAILURE";
    case CL_INVALID_VALUE: return "CL_INVALID_VALUE";
    case CL_INVALID_DEVICE: return "CL_INVALID_DEVICE";
    case CL_INVALID_CONTEXT: return "CL_INVALID_CONTEXT";
    case CL_INVALID_QUEUE_PROPERTIES: return "CL_INVALID_QUEUE_PROPERTIES";
    case CL_INVALID_COMMAND_QUEUE: return "CL_INVALID_COMMAND_QUEUE";
    case CL_INVALID_MEM_OBJECT: return "CL_INVALID_MEM_OBJECT";
    case CL_INVALID_BUFFER_SIZE: return "CL_INVALID_BUFFER_SIZE";
    case CL_INVALID_BUILD_OPTIONS: return "CL_INVALID_BUILD_OPTIONS";
    case CL_INVALID_PROGRAM: return "CL_INVALID_PROGRAM";
    case CL_INVALID_PROGRAM_EXECUTABLE: return "CL_INVALID_PROGRAM_EXECUTABLE";
    case CL_INVALID_KERNEL_NAME: return "CL_INVALID_KERNEL_NAME";
    case CL_INVALID_KERNEL: return "CL_INVALID_KERNEL";
    case CL_INVALID_ARG_INDEX: return "CL_INVALID_ARG_INDEX";
    case CL_INVALID_ARG_VALUE: return "CL_INVALID_ARG_VALUE";
    case CL_INVALID_ARG_SIZE: return "CL_INVALID_ARG_SIZE";
    case CL_INVALID_KERNEL_ARGS: return "CL_INVALID_KERNEL_ARGS";
    case CL_INVALID_WORK_DIMENSION: return "CL_INVALID_WORK_DIMENSION";
    case CL_INVALID_WORK_GROUP_SIZE: return "CL_INVALID_WORK_GROUP_SIZE";
    case CL_INVALID_WORK_ITEM_SIZE: return "CL_INVALID_WORK_ITEM_SIZE";
    case CL_INVALID_GLOBAL_OFFSET: return "CL_INVALID_GLOBAL_OFFSET";
    case CL_INVALID_EVENT: return "CL_INVALID_EVENT";
    case CL_INVALID_OPERATION: return "CL_INVALID_OPERATION";
    case CL_INVALID_GLOBAL_WORK_SIZE: return "CL_INVALID_GLOBAL_WORK_SIZE";
    case CL_INVALID_PLATFORM: return "CL_INVALID_PLATFORM";
    default: return "(other)";
    }
}

static void clFail(cl_int e, const char* what, int line) {
    fprintf(stderr, "OpenCL error: %s (%d)\n  at host.c:%d\n  in %s\n", clErrName(e), (int)e, line, what);
    exit(2);
}
#define CL_CHECK(call) do { cl_int err_ = (call); if (err_ != CL_SUCCESS) clFail(err_, #call, __LINE__); } while (0)
#define CL_CHECK_ERR(err_, what) do { if ((err_) != CL_SUCCESS) clFail((err_), what, __LINE__); } while (0)

static double wallMs(void) {
#ifdef _WIN32
    LARGE_INTEGER f, c;
    QueryPerformanceFrequency(&f);
    QueryPerformanceCounter(&c);
    return (double)c.QuadPart * 1000.0 / (double)f.QuadPart;
#else
    struct timespec ts;
    clock_gettime(CLOCK_MONOTONIC, &ts);
    return (double)ts.tv_sec * 1000.0 + (double)ts.tv_nsec / 1e6;
#endif
}

// Event profiling. A runtime that cannot report timestamps (CL_PROFILING_INFO_NOT_AVAILABLE) gives -1 and the
// harness switches the rate to wall time instead of stopping; the count of such events is reported.
static int gProfilingFailures = 0;
static double eventMs(cl_event e) {
    cl_ulong t0 = 0, t1 = 0;
    if (clGetEventProfilingInfo(e, CL_PROFILING_COMMAND_START, sizeof(t0), &t0, NULL) != CL_SUCCESS ||
        clGetEventProfilingInfo(e, CL_PROFILING_COMMAND_END, sizeof(t1), &t1, NULL) != CL_SUCCESS) { ++gProfilingFailures; return -1.0; }
    return (double)(t1 - t0) / 1e6;
}

static double spanMs(cl_event first, cl_event last) {
    cl_ulong t0 = 0, t1 = 0;
    if (clGetEventProfilingInfo(first, CL_PROFILING_COMMAND_START, sizeof(t0), &t0, NULL) != CL_SUCCESS ||
        clGetEventProfilingInfo(last, CL_PROFILING_COMMAND_END, sizeof(t1), &t1, NULL) != CL_SUCCESS) { ++gProfilingFailures; return -1.0; }
    return (double)(t1 - t0) / 1e6;
}

static void* loadSym(const char* name) {
#ifdef _WIN32
    HMODULE m = GetModuleHandleA("OpenCL.dll");
    return m ? (void*)GetProcAddress(m, name) : NULL;
#else
    return dlsym(RTLD_DEFAULT, name);
#endif
}

// ---------------------------------------------------------------------------------------------
// Host reference

static const uint32_t SEEDW[8] = IGNEUM_SEEDW_INIT;

#if IGNEUM_DATASET_MODE == 0
// Same closed form as ds_elem in kernel.cl and datasetElem in proto-metal/main.swift.
static uint32_t host_ds_elem(uint32_t i, uint32_t d0, uint32_t d1) {
    uint32_t x = i ^ d0;
    x *= 0x9E3779B1u; x ^= x >> 15;
    x += d1;
    x *= 0x85EBCA77u; x ^= x >> 13;
    x *= 0xC2B2AE3Du; x ^= x >> 16;
    return x;
}
#else
static uint64_t fnv1a64(const void* p, size_t n) {
    const uint8_t* b = (const uint8_t*)p;
    uint64_t h = 0xcbf29ce484222325ull;
    size_t i;
    for (i = 0; i < n; ++i) { h ^= b[i]; h *= 0x100000001b3ull; }
    return h;
}
static const uint32_t CACHE_WORDS_HOST = 1u << IGNEUM_CACHE_LOG2_WORDS;
static uint32_t* hCache = NULL;
/* dataset[w] through the pack's own mh_word (memhard.h), which carries the pack's item-to-word layout (era layout,
 * 5 October 2026; the former w >> 4 / w & 15 here failed the random points of every interleaved pack). */
static uint32_t host_ds_word(uint32_t w) {
    return mh_word(hCache, w);
}
#endif

// ---------------------------------------------------------------------------------------------
// Options

typedef struct {
    int datasetMib;
    int batchLog2;
    int batches;
    int groupWarps;
    int sweep;
    int device;          // flat index into the enumerated list, -1 = first GPU
    int exchange;        // 0 auto, 1 force local-memory fallback, 2 force sub-group shuffles
    int list;
    int timeWall;        // 1 = rate from wall time, 0 = from device event profiling, -1 = auto (wall on the Apple platform)
    const char* kernelPath;
    const char* extraOpts;
    int serve;           // --serve: GPU worker for igneum-miner --worker (jobs on stdin), 3 October 2026
    int noPrepare;       // --no-prepare: serve without the prepare command (ready line says "prepare 0"), to test the miner's fallback
    int kernelGiven;     // --kernel was passed
    const char* vendor;  // --vendor S: pick the first GPU whose vendor string contains S (default: first GPU of any vendor)
    const char* packDir; // --pack D (serve only): generic mode, the pack is read from D at run time; the compiled-in pack is
                         // then only the build-time placeholder of the prebuilt exe (4 October 2026)
    int readback;        // --readback: 0 select (a GPU-side pass reads back only the hits and 34 sentinel words), 1 full
                         // (every output word comes back, 8 bytes per nonce, the path before 5 October 2026)
    int memprobe;        // --memprobe: dependent-load latency and throughput, independent-load throughput and an ALU
                         // chain on the chosen device, no pack needed (5 October 2026, the 9070 XT on the eGPU)
    int probeMib;        // --probe-mib N: --memprobe at that one buffer size only (default 0 = 4, 64 and 1024 MiB)
    int benchPack;       // --bench-pack: with --pack D, build and self-test the pack at run time (as --serve does) and time
                         // igneum_hash_bound with the pack's seed words as init words; read-width experiment, 5 October 2026
    int warps;           // --warps N: persistent warps for a variant-5 pack (IGNEUM_PERSISTENT_WARPS); 0 = 2048
} Options;

static int packMib(void) { return (int)(((1ull << IGNEUM_DATASET_LOG2) * 4ull) >> 20); }

static void usage(void) {
    printf(
        "igneum-bench-cl [--list] [--device D] [--dataset-mib N] [--sweep] [--batch-log2 24] [--batches 5] [--group-warps 1]\n"
        "                [--exchange auto|local|subgroup] [--kernel path/to/kernel.cl] [--build-opts \"...\"]\n"
        "  --list            print every OpenCL platform and device, then exit\n"
        "  --device D        device index from the list (default: the first GPU, else device 0)\n"
        "  --dataset-mib N   dataset size in MiB, power of two (default 1024; vectors are only checked at %d MiB)\n"
        "  --sweep           run 4, 64, 256, 512 and 1024 MiB in sequence (same sweep as the Mac and the CUDA harness)\n"
        "  --batch-log2 B    nonces per batch = 2^B (default 24)\n"
        "  --batches N       timed batches after one warm-up batch (default 5)\n"
        "  --group-warps W   32-lane units per work-group, 1..8 (default 1 = one work-group per unit; sub-group shuffles need 1)\n"
        "  --exchange M      auto (default): sub-group shuffles when the device has them and its sub-group size is 32, else local memory\n"
        "                    local: force the local-memory exchange; subgroup: require sub-group shuffles or fail\n"
        "  --kernel P        path to the pack's kernel.cl (default: the path compiled in, %s)\n"
        "  --build-opts S    extra options appended to clBuildProgram (for example \"-cl-std=CL2.0\")\n"
        "  --time T          event (default): hashes/s from device event profiling, like cudaEvent time; wall: from host wall time.\n"
        "                    Apple's OpenCL runtime reports unusable event timestamps, so wall is the default on the Apple platform.\n"
        "  --vendor S        choose the first GPU whose vendor string contains S (for example \"Advanced Micro Devices\"); fails if none\n"
        "  --serve           GPU worker for igneum-miner --worker: reads \"job ...\" lines on stdin, prints found/done lines.\n"
        "  --no-prepare      with --serve: no prepare support (the miner then falls back to exit 42 at a seed change).\n"
        "                    Builds the pack's kernel_bound.cl (next to the compiled-in kernel.cl) unless --kernel says otherwise.\n"
        "  --pack D          with --serve: serve the pack in directory D (program.h, seeds.txt, vectors.h, kernel_bound.cl), whatever\n"
        "                    pack this exe was built against; it is self-tested against its vectors.h first (the one-click worker)\n"
        "  --readback M      with --serve: select (default) reads back only the hits and 34 sentinel words of each dispatch through a\n"
        "                    GPU-side pass; full reads back every output (8 bytes per nonce). IGNEUM_READBACK=full does the same.\n"
        "  --bench-pack      with --pack D: read the pack at run time, build and self-test it, time its bound kernel (one exe, any pack)\n"
        "  --warps N         persistent warps for a variant-5 pack (a 1 MiB scratch each; default 2048; the batch rounds to 32 x N)\n"
        "  --memprobe        no pack: dependent random loads (latency and throughput against lanes in flight), independent random\n"
        "                    loads and an ALU chain on the chosen device, at 4, 64 and 1024 MiB (--probe-mib N for one size)\n", packMib(), IGNEUM_KERNEL_PATH);
}

static int isPow2(long long v) { return v > 0 && (v & (v - 1)) == 0; }
static int log2u32(uint32_t v) { int n = 0; while (v > 1u) { v >>= 1; ++n; } return n; }

static Options parseArgs(int argc, char** argv) {
    Options o;
    int i;
    o.datasetMib = 1024; o.batchLog2 = 24; o.batches = 5; o.groupWarps = 1; o.sweep = 0; o.device = -1;
    o.exchange = 0; o.list = 0; o.timeWall = -1; o.kernelPath = IGNEUM_KERNEL_PATH; o.extraOpts = ""; o.serve = 0; o.noPrepare = 0; o.kernelGiven = 0; o.vendor = NULL; o.packDir = NULL;
    o.readback = (getenv("IGNEUM_READBACK") && strcmp(getenv("IGNEUM_READBACK"), "full") == 0) ? 1 : 0; o.memprobe = 0; o.probeMib = 0; o.benchPack = 0; o.warps = 0;
    for (i = 1; i < argc; ++i) {
        const char* a = argv[i];
        int needs = (strcmp(a, "--dataset-mib") == 0 || strcmp(a, "--batch-log2") == 0 || strcmp(a, "--batches") == 0 ||
                     strcmp(a, "--group-warps") == 0 || strcmp(a, "--device") == 0 || strcmp(a, "--exchange") == 0 ||
                     strcmp(a, "--kernel") == 0 || strcmp(a, "--build-opts") == 0 || strcmp(a, "--time") == 0);
        if (needs && i + 1 >= argc) { usage(); exit(2); }
        if (strcmp(a, "--dataset-mib") == 0) o.datasetMib = atoi(argv[++i]);
        else if (strcmp(a, "--batch-log2") == 0) o.batchLog2 = atoi(argv[++i]);
        else if (strcmp(a, "--batches") == 0) o.batches = atoi(argv[++i]);
        else if (strcmp(a, "--group-warps") == 0) o.groupWarps = atoi(argv[++i]);
        else if (strcmp(a, "--device") == 0) o.device = atoi(argv[++i]);
        else if (strcmp(a, "--kernel") == 0) { o.kernelPath = argv[++i]; o.kernelGiven = 1; }
        else if (strcmp(a, "--serve") == 0) o.serve = 1;
        else if (strcmp(a, "--no-prepare") == 0) o.noPrepare = 1;
        else if (strcmp(a, "--vendor") == 0) { if (i + 1 >= argc) { usage(); exit(2); } o.vendor = argv[++i]; }
        else if (strcmp(a, "--pack") == 0) { if (i + 1 >= argc) { usage(); exit(2); } o.packDir = argv[++i]; }
        else if (strcmp(a, "--readback") == 0) {
            const char* m;
            if (i + 1 >= argc) { usage(); exit(2); }
            m = argv[++i];
            if (strcmp(m, "select") == 0) o.readback = 0;
            else if (strcmp(m, "full") == 0) o.readback = 1;
            else { printf("--readback must be select or full\n"); exit(2); }
        }
        else if (strcmp(a, "--memprobe") == 0) o.memprobe = 1;
        else if (strcmp(a, "--bench-pack") == 0) o.benchPack = 1;
        else if (strcmp(a, "--warps") == 0) { if (i + 1 >= argc) { usage(); exit(2); } o.warps = atoi(argv[++i]); }
        else if (strcmp(a, "--probe-mib") == 0) { if (i + 1 >= argc) { usage(); exit(2); } o.probeMib = atoi(argv[++i]); }
        else if (strcmp(a, "--build-opts") == 0) o.extraOpts = argv[++i];
        else if (strcmp(a, "--time") == 0) {
            const char* m = argv[++i];
            if (strcmp(m, "event") == 0) o.timeWall = 0;
            else if (strcmp(m, "wall") == 0) o.timeWall = 1;
            else { printf("--time must be event or wall\n"); exit(2); }
        }
        else if (strcmp(a, "--exchange") == 0) {
            const char* m = argv[++i];
            if (strcmp(m, "auto") == 0) o.exchange = 0;
            else if (strcmp(m, "local") == 0) o.exchange = 1;
            else if (strcmp(m, "subgroup") == 0) o.exchange = 2;
            else { printf("--exchange must be auto, local or subgroup\n"); exit(2); }
        }
        else if (strcmp(a, "--sweep") == 0) o.sweep = 1;
        else if (strcmp(a, "--list") == 0) o.list = 1;
        else if (strcmp(a, "-h") == 0 || strcmp(a, "--help") == 0) { usage(); exit(0); }
        else { printf("unknown argument %s\n", a); usage(); exit(2); }
    }
    if (!isPow2(o.datasetMib) || o.datasetMib < 1 || o.datasetMib > 16384) { printf("--dataset-mib must be a power of two between 1 and 16384\n"); exit(2); }
    if (o.batchLog2 < 10 || o.batchLog2 > 28) { printf("--batch-log2 must be between 10 and 28\n"); exit(2); }
    if (o.batches < 1) { printf("--batches must be at least 1\n"); exit(2); }
    if (o.groupWarps < 1 || o.groupWarps > 8) { printf("--group-warps must be between 1 and 8\n"); exit(2); }
    return o;
}

// ---------------------------------------------------------------------------------------------
// Devices

typedef struct {
    cl_platform_id platform;
    cl_device_id device;
    char platformName[256], platformVersion[256];
    char name[256], vendor[256], version[256], driver[256], cVersion[256];
    char* extensions;
    cl_device_type type;
    cl_uint computeUnits, clockMHz;
    cl_ulong globalMem, maxAlloc, localMem;
    size_t maxWorkGroup;
    int cMajor, cMinor;          // OpenCL C version
    int dMajor, dMinor;          // device (platform profile) version
    cl_uint amdWavefront, nvWarp; // 0 if not reported
    int dupOf;                   // index of the same card on a newer platform of the same vendor, -1 if none (5 October 2026)
    char pci[32];                 // "01:00.0" (bus:device.function) when the vendor extension reports it, else ""
} DeviceInfo;

/* The driver version as a number for ordering ("3683.0 (PAL,LC)" -> 3683.0; "617.14" -> 617.14; 0 when unreadable). */
static double driverNumber(const char* driver) {
    const char* p = driver;
    while (*p && (*p < '0' || *p > '9')) ++p;
    return *p ? strtod(p, NULL) : 0.0;
}

/* Two AMD platforms are registered after a driver update on Windows (PC 1, 5 October 2026: 32.0.21042 and 32.0.32015,
 * OpenCL driver strings 3652.0 and 3683.0): every card is listed twice, the app ran two workers on one 9070 XT, and
 * each got half. The same card on the same-named platform with a different platform version is the one card; the
 * entry whose driver number is lower is the duplicate. Two real cards of one model sit on the SAME platform and are
 * never folded. Indices stay flat (the app passes them back as --device). Returns how many duplicates were marked. */
static int markDuplicates(DeviceInfo* list, int n) {
    int i, j, marked = 0;
    for (i = 0; i < n; ++i) list[i].dupOf = -1;
    for (i = 0; i < n; ++i) {
        if (list[i].dupOf >= 0) continue;
        for (j = i + 1; j < n; ++j) {
            int older;
            if (list[j].dupOf >= 0) continue;
            if (strcmp(list[i].platformName, list[j].platformName) != 0) continue;
            if (strcmp(list[i].platformVersion, list[j].platformVersion) == 0) continue;
            if (strcmp(list[i].name, list[j].name) != 0 || strcmp(list[i].vendor, list[j].vendor) != 0) continue;
            if (list[i].globalMem != list[j].globalMem || list[i].computeUnits != list[j].computeUnits) continue;
            older = driverNumber(list[j].driver) < driverNumber(list[i].driver) ? j : i;
            if (older == j) { list[j].dupOf = i; }
            else { list[i].dupOf = j; }
            ++marked;
            if (older == i) break;   /* i itself is the duplicate; j stays the real one */
        }
    }
    /* three registrations of one card: every duplicate points at the one that stays, not at another duplicate */
    for (i = 0; i < n; ++i) {
        int k = list[i].dupOf, hops = 0;
        while (k >= 0 && list[k].dupOf >= 0 && hops++ < n) k = list[k].dupOf;
        if (list[i].dupOf >= 0) list[i].dupOf = k;
    }
    return marked;
}

static void devStr(cl_device_id d, cl_device_info what, char* out, size_t n) {
    out[0] = 0;
    clGetDeviceInfo(d, what, n - 1, out, NULL);
    out[n - 1] = 0;
}

static const char* typeName(cl_device_type t) {
    if (t & CL_DEVICE_TYPE_GPU) return "GPU";
    if (t & CL_DEVICE_TYPE_CPU) return "CPU";
    if (t & CL_DEVICE_TYPE_ACCELERATOR) return "accelerator";
    return "other";
}

static int enumerateDevices(DeviceInfo** outList) {
    cl_uint np = 0, p;
    cl_platform_id plats[16];
    DeviceInfo* list = NULL;
    int n = 0;
    cl_int e = clGetPlatformIDs(16, plats, &np);
    if (e != CL_SUCCESS || np == 0) { *outList = NULL; return 0; }
    for (p = 0; p < np; ++p) {
        cl_uint nd = 0, d;
        cl_device_id devs[32];
        char pname[256] = {0}, pver[256] = {0};
        clGetPlatformInfo(plats[p], CL_PLATFORM_NAME, sizeof(pname) - 1, pname, NULL);
        clGetPlatformInfo(plats[p], CL_PLATFORM_VERSION, sizeof(pver) - 1, pver, NULL);
        if (clGetDeviceIDs(plats[p], CL_DEVICE_TYPE_ALL, 32, devs, &nd) != CL_SUCCESS) continue;
        for (d = 0; d < nd; ++d) {
            DeviceInfo di;
            size_t extLen = 0;
            memset(&di, 0, sizeof(di));
            di.platform = plats[p]; di.device = devs[d];
            strncpy(di.platformName, pname, 255); strncpy(di.platformVersion, pver, 255);
            devStr(devs[d], CL_DEVICE_NAME, di.name, sizeof(di.name));
            devStr(devs[d], CL_DEVICE_VENDOR, di.vendor, sizeof(di.vendor));
            devStr(devs[d], CL_DEVICE_VERSION, di.version, sizeof(di.version));
            devStr(devs[d], CL_DRIVER_VERSION, di.driver, sizeof(di.driver));
            devStr(devs[d], CL_DEVICE_OPENCL_C_VERSION, di.cVersion, sizeof(di.cVersion));
            clGetDeviceInfo(devs[d], CL_DEVICE_EXTENSIONS, 0, NULL, &extLen);
            di.extensions = (char*)calloc(extLen + 1, 1);
            if (extLen) clGetDeviceInfo(devs[d], CL_DEVICE_EXTENSIONS, extLen, di.extensions, NULL);
            clGetDeviceInfo(devs[d], CL_DEVICE_TYPE, sizeof(di.type), &di.type, NULL);
            clGetDeviceInfo(devs[d], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(di.computeUnits), &di.computeUnits, NULL);
            clGetDeviceInfo(devs[d], CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof(di.clockMHz), &di.clockMHz, NULL);
            clGetDeviceInfo(devs[d], CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(di.globalMem), &di.globalMem, NULL);
            clGetDeviceInfo(devs[d], CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(di.maxAlloc), &di.maxAlloc, NULL);
            clGetDeviceInfo(devs[d], CL_DEVICE_LOCAL_MEM_SIZE, sizeof(di.localMem), &di.localMem, NULL);
            clGetDeviceInfo(devs[d], CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(di.maxWorkGroup), &di.maxWorkGroup, NULL);
            if (sscanf(di.cVersion, "OpenCL C %d.%d", &di.cMajor, &di.cMinor) != 2) { di.cMajor = 1; di.cMinor = 2; }
            if (sscanf(di.version, "OpenCL %d.%d", &di.dMajor, &di.dMinor) != 2) { di.dMajor = 1; di.dMinor = 2; }
            if (strstr(di.extensions, "cl_amd_device_attribute_query")) {
                ig_topology_amd topo;
                memset(&topo, 0, sizeof(topo));
                clGetDeviceInfo(devs[d], IG_CL_DEVICE_WAVEFRONT_WIDTH_AMD, sizeof(di.amdWavefront), &di.amdWavefront, NULL);
                if (clGetDeviceInfo(devs[d], IG_CL_DEVICE_TOPOLOGY_AMD, sizeof(topo), &topo, NULL) == CL_SUCCESS && topo.raw.type == IG_CL_DEVICE_TOPOLOGY_TYPE_PCIE_AMD)
                    snprintf(di.pci, sizeof(di.pci), "%02x:%02x.%x", (unsigned)(unsigned char)topo.pcie.bus, (unsigned)(unsigned char)topo.pcie.device, (unsigned)(unsigned char)topo.pcie.function);
            }
            if (strstr(di.extensions, "cl_nv_device_attribute_query")) {
                cl_uint bus = 0, slot = 0;
                clGetDeviceInfo(devs[d], IG_CL_DEVICE_WARP_SIZE_NV, sizeof(di.nvWarp), &di.nvWarp, NULL);
                if (clGetDeviceInfo(devs[d], IG_CL_DEVICE_PCI_BUS_ID_NV, sizeof(bus), &bus, NULL) == CL_SUCCESS && clGetDeviceInfo(devs[d], IG_CL_DEVICE_PCI_SLOT_ID_NV, sizeof(slot), &slot, NULL) == CL_SUCCESS)
                    snprintf(di.pci, sizeof(di.pci), "%02x:%02x.0", (unsigned)(bus & 0xff), (unsigned)(slot & 0xff));
            }
            list = (DeviceInfo*)realloc(list, sizeof(DeviceInfo) * (size_t)(n + 1));
            list[n++] = di;
        }
    }
    if (n) markDuplicates(list, n);
    *outList = list;
    return n;
}

static void printDevice(int idx, const DeviceInfo* d, int chosen) {
    const char* subExt = strstr(d->extensions, "cl_khr_subgroup_shuffle") ? "cl_khr_subgroup_shuffle" :
                         strstr(d->extensions, "cl_intel_subgroups") ? "cl_intel_subgroups" :
                         strstr(d->extensions, "cl_khr_subgroups") ? "cl_khr_subgroups (no shuffle extension)" : "none";
    if (d->dupOf >= 0) {
        /* Hidden from the app's card list: its parser takes only lines that start with "[" (detect.rs). The index
         * is still valid for --device, so a run on the older platform stays possible for a comparison. */
        printf(" dup [%d] %s  |  %s (%s): the same card as [%d] on an older platform (driver %s); hidden, use [%d]\n",
               idx, d->name, d->platformName, d->platformVersion, d->dupOf, d->driver, d->dupOf);
        return;
    }
    printf("%s[%d] %s  |  %s (%s)\n", chosen ? "*" : " ", idx, d->name, d->platformName, d->platformVersion);
    // the app's detect.rs reads this line: type, vendor, driver, the compute units, and the PCI address when known
    printf("      %s, vendor %s, driver %s, %s, %u compute units, %u MHz", typeName(d->type), d->vendor, d->driver, d->cVersion, d->computeUnits, d->clockMHz);
    if (d->pci[0]) printf(", pci %s", d->pci);
    printf("\n");
    printf("      global %llu MiB, max alloc %llu MiB, local %llu KiB, max work-group %llu, sub-group extension: %s",
           (unsigned long long)(d->globalMem >> 20), (unsigned long long)(d->maxAlloc >> 20), (unsigned long long)(d->localMem >> 10),
           (unsigned long long)d->maxWorkGroup, subExt);
    if (d->amdWavefront) printf(", AMD wavefront width %u", d->amdWavefront);
    if (d->nvWarp) printf(", NVIDIA warp size %u", d->nvWarp);
    printf("\n");
}

// ---------------------------------------------------------------------------------------------
// Program build

static char* readFile(const char* path, size_t* len) {
    FILE* f = fopen(path, "rb");
    char* buf;
    long n;
    if (!f) return NULL;
    fseek(f, 0, SEEK_END); n = ftell(f); fseek(f, 0, SEEK_SET);
    if (n < 0) { fclose(f); return NULL; }
    buf = (char*)malloc((size_t)n + 1);
    if (fread(buf, 1, (size_t)n, f) != (size_t)n) { fclose(f); free(buf); return NULL; }
    buf[n] = 0;
    fclose(f);
    *len = (size_t)n;
    return buf;
}

typedef struct {
    cl_context ctx;
    cl_command_queue q;
    cl_program prog;
    cl_kernel kHash, kCacheFill, kBuild, kFill;
    cl_kernel kHashBound;     // igneum_hash_bound (serve mode; NULL when the source has none)
    int exchange;             // 0 local memory, 1 khr sub-group shuffle, 2 intel
    size_t subGroupSize;      // as queried for a 32-item work-group, 0 if not queried
    char exchangeNote[512];
    char buildOptions[512];
    int groupSize;            // work-group size the program was built for (IGNEUM_GROUP = 32 x group-warps)
} Device;

static const char* exchangeName(int m) { return m == 1 ? "sub_group_shuffle_xor (cl_khr_subgroup_shuffle)" : m == 2 ? "intel_sub_group_shuffle_xor (cl_intel_subgroups)" : "local-memory exchange with barrier"; }

// Returns 0 on success, 1 on build failure (log printed).
static int buildProgram(Device* dv, const DeviceInfo* di, const char* src, size_t srcLen, int exchangeMode, int groupSize, const char* extra) {
    cl_int err = 0;
    double tb = wallMs();   /* the pack's compile cost (Counter ASIC 3.0 item 2, 6 October 2026): printed as one line below */
    const char* std;
    // The sub-group built-ins need OpenCL C 2.0 or 3.0. OpenCL 3.0 devices may report "OpenCL C 1.2" as the default
    // CL_DEVICE_OPENCL_C_VERSION while supporting 3.0 (the 3.0 API lists all versions; the 1.2 API cannot ask), so
    // the device version counts as well. The local-memory variant is always built as OpenCL C 1.2, the same text everywhere.
    int major = di->cMajor > di->dMajor ? di->cMajor : di->dMajor;
    if (exchangeMode == 0) std = "-cl-std=CL1.2";
    else if (major >= 3) std = "-cl-std=CL3.0";
    else if (major >= 2) std = "-cl-std=CL2.0";
    else std = "-cl-std=CL1.2";
    snprintf(dv->buildOptions, sizeof(dv->buildOptions), "%s -D IGNEUM_GROUP=%d -D IGNEUM_EXCHANGE=%d %s", std, groupSize, exchangeMode, extra);
    dv->prog = clCreateProgramWithSource(dv->ctx, 1, &src, &srcLen, &err);
    CL_CHECK_ERR(err, "clCreateProgramWithSource");
    err = clBuildProgram(dv->prog, 1, &di->device, dv->buildOptions, NULL, NULL);
    if (err != CL_SUCCESS) {
        size_t logLen = 0;
        char* log;
        clGetProgramBuildInfo(dv->prog, di->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &logLen);
        log = (char*)calloc(logLen + 1, 1);
        if (logLen) clGetProgramBuildInfo(dv->prog, di->device, CL_PROGRAM_BUILD_LOG, logLen, log, NULL);
        printf("build FAILED (%s) with options \"%s\"\n--- build log ---\n%s\n--- end of build log ---\n", clErrName(err), dv->buildOptions, log);
        free(log);
        clReleaseProgram(dv->prog); dv->prog = NULL;
        return 1;
    }
    /* The OpenCL build time of this pack's kernel text (clBuildProgram alone), the equivalent of the CUDA worker's
     * NVRTC line: a per-day item-derivation program (item 2) sits inside every hash-kernel compile, so its cost is
     * read here. Wall time, printed before the kernels are created. */
    printf("build %.1f ms clBuildProgram (exchange %d, group %d)\n", wallMs() - tb, exchangeMode, groupSize);
    dv->kHash = clCreateKernel(dv->prog, "igneum_hash", &err); CL_CHECK_ERR(err, "clCreateKernel igneum_hash");
    dv->kHashBound = clCreateKernel(dv->prog, "igneum_hash_bound", &err);
    if (err != CL_SUCCESS) dv->kHashBound = NULL;   /* kernel.cl without the bound kernel: fine outside --serve */
#if IGNEUM_DATASET_MODE == 1
    dv->kCacheFill = clCreateKernel(dv->prog, "igneum_cache_fill", &err); CL_CHECK_ERR(err, "clCreateKernel igneum_cache_fill");
    dv->kBuild = clCreateKernel(dv->prog, "igneum_build", &err); CL_CHECK_ERR(err, "clCreateKernel igneum_build");
#else
    dv->kFill = clCreateKernel(dv->prog, "igneum_fill", &err); CL_CHECK_ERR(err, "clCreateKernel igneum_fill");
#endif
    return 0;
}

static void releaseProgram(Device* dv) {
    if (dv->kHash) clReleaseKernel(dv->kHash);
    if (dv->kHashBound) clReleaseKernel(dv->kHashBound);
    if (dv->kCacheFill) clReleaseKernel(dv->kCacheFill);
    if (dv->kBuild) clReleaseKernel(dv->kBuild);
    if (dv->kFill) clReleaseKernel(dv->kFill);
    if (dv->prog) clReleaseProgram(dv->prog);
    dv->kHash = dv->kCacheFill = dv->kBuild = dv->kFill = dv->kHashBound = NULL; dv->prog = NULL;
}

// Sub-group size of kernel k for a work-group of `local` items. 0 if the query is unavailable (reason in *why).
static size_t querySubGroupSizeOf(cl_kernel k, const DeviceInfo* di, size_t local, char* why, size_t whyLen) {
    ig_pfn_subgroup_info fn = (ig_pfn_subgroup_info)clGetExtensionFunctionAddressForPlatform(di->platform, "clGetKernelSubGroupInfoKHR");
    const char* via = "clGetKernelSubGroupInfoKHR";
    size_t sg = 0;
    cl_int e;
    if (!fn) { fn = (ig_pfn_subgroup_info)loadSym("clGetKernelSubGroupInfo"); via = "clGetKernelSubGroupInfo (OpenCL 2.1 core, through the loader)"; }
    if (!fn) { snprintf(why, whyLen, "the sub-group size could not be queried (neither clGetKernelSubGroupInfoKHR nor clGetKernelSubGroupInfo is available)"); return 0; }
    e = fn(k, di->device, IG_CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE, sizeof(local), &local, sizeof(sg), &sg, NULL);
    if (e != CL_SUCCESS) { snprintf(why, whyLen, "the sub-group size query failed: %s returned %s (%d)", via, clErrName(e), (int)e); return 0; }
    snprintf(why, whyLen, "queried through %s", via);
    return sg;
}
static size_t querySubGroupSize(const Device* dv, const DeviceInfo* di, size_t local, char* why, size_t whyLen) {
    return querySubGroupSizeOf(dv->kHash, di, local, why, whyLen);
}

/* The compiled kernel as the driver sees it, on every exchange path (5 October 2026, the 9070 XT question): the
 * work-group limit, the preferred multiple (the wave width the compiler chose), local memory, private memory (scratch:
 * anything above 0 means spilled registers, which on AMD costs a memory round trip per spill), and the sub-group size
 * for the built work-group size (wave32 or wave64 on RDNA; the local-memory path never queried it before). */
#define IG_CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE 0x11B3
#define IG_CL_KERNEL_PRIVATE_MEM_SIZE 0x11B4
static void printKernelInfo(const DeviceInfo* di, cl_kernel k, const char* kernelName, int groupSize, const char* prefix) {
    size_t wg = 0, mult = 0;
    cl_ulong lmem = 0, pmem = 0;
    char why[256];
    size_t sg;
    clGetKernelWorkGroupInfo(k, di->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(wg), &wg, NULL);
    clGetKernelWorkGroupInfo(k, di->device, IG_CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, sizeof(mult), &mult, NULL);
    clGetKernelWorkGroupInfo(k, di->device, CL_KERNEL_LOCAL_MEM_SIZE, sizeof(lmem), &lmem, NULL);
    clGetKernelWorkGroupInfo(k, di->device, IG_CL_KERNEL_PRIVATE_MEM_SIZE, sizeof(pmem), &pmem, NULL);
    sg = querySubGroupSizeOf(k, di, (size_t)groupSize, why, sizeof(why));
    printf("%skernel: %s max work-group %llu, preferred multiple %llu, local memory %llu bytes, private memory %llu bytes%s, work-group %d, sub-group size %llu (%s)%s\n",
           prefix, kernelName, (unsigned long long)wg, (unsigned long long)mult, (unsigned long long)lmem, (unsigned long long)pmem,
           pmem ? " (SPILLED: registers in scratch memory)" : "", groupSize, (unsigned long long)sg, why,
           (sg > (size_t)groupSize) ? " (the work-group fills only part of a wave: see --group-warps)" : "");
    fflush(stdout);
}

// Decide the exchange implementation and build. See WAVEFRONT.md for the rule.
static void setupProgram(Device* dv, const DeviceInfo* di, const Options* o, const char* src, size_t srcLen) {
    int groupSize = 32 * o->groupWarps;
    int want = 0;
    dv->groupSize = groupSize;
    const char* why = "";
    if (strstr(di->extensions, "cl_khr_subgroup_shuffle")) want = 1;
    else if (strstr(di->extensions, "cl_intel_subgroups")) want = 2;
    else why = "device lists no sub-group shuffle extension";
    if (o->exchange == 1) { want = 0; why = "forced by --exchange local"; }
    if (want != 0 && o->groupWarps != 1) {
        if (o->exchange == 2) { printf("FAIL: --exchange subgroup needs --group-warps 1 (one work-group = one 32-lane unit)\n"); exit(2); }
        want = 0; why = "--group-warps is not 1, so a work-group is not one 32-lane unit";
    }
    if (want == 0 && o->exchange == 2) { printf("FAIL: --exchange subgroup requested but %s\n", why); exit(2); }

    if (want != 0) {
        if (buildProgram(dv, di, src, srcLen, want, groupSize, o->extraOpts) != 0) {
            if (o->exchange == 2) { printf("FAIL: the sub-group variant did not compile\n"); exit(2); }
            want = 0; why = "the sub-group variant did not compile (log above)";
        } else {
            static char qwhy[256];
            size_t sg = querySubGroupSize(dv, di, 32, qwhy, sizeof(qwhy));
            if (sg == 0) {
                // Second choice: a probe kernel from the same build reports get_sub_group_size() for a 32-item work-group.
                // Weaker than the per-kernel query (a compiler may pick the wave width per kernel), and the note says so.
                cl_int perr = 0;
                cl_kernel kp = clCreateKernel(dv->prog, "igneum_probe_subgroup", &perr);
                if (perr == CL_SUCCESS) {
                    cl_uint probe[2] = { 0u, 0u };
                    cl_mem pb = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, sizeof(probe), NULL, &perr);
                    size_t g = 32, l = 32;
                    if (perr == CL_SUCCESS && clSetKernelArg(kp, 0, sizeof(cl_mem), &pb) == CL_SUCCESS &&
                        clEnqueueNDRangeKernel(dv->q, kp, 1, NULL, &g, &l, 0, NULL, NULL) == CL_SUCCESS &&
                        clEnqueueReadBuffer(dv->q, pb, CL_TRUE, 0, sizeof(probe), probe, 0, NULL, NULL) == CL_SUCCESS && probe[0] != 0u) {
                        static char pwhy[400];
                        snprintf(pwhy, sizeof(pwhy), "%s; probe kernel reports get_sub_group_size() %u and %u sub-group(s) for a 32-item work-group (weaker than the per-kernel query)", qwhy, probe[0], probe[1]);
                        sg = probe[0];
                        strncpy(qwhy, pwhy, sizeof(qwhy) - 1); qwhy[sizeof(qwhy) - 1] = 0;
                    }
                    if (pb) clReleaseMemObject(pb);
                    clReleaseKernel(kp);
                }
            }
            dv->subGroupSize = sg;
            if (sg == 32) {
                dv->exchange = want;
                snprintf(dv->exchangeNote, sizeof(dv->exchangeNote), "%s, sub-group size %llu for a 32-item work-group (%s)", exchangeName(want), (unsigned long long)sg, qwhy);
                return;
            }
            releaseProgram(dv);
            if (sg == 0) why = qwhy;
            else why = "the sub-group size for a 32-item work-group is not 32";
            if (o->exchange == 2) { printf("FAIL: --exchange subgroup requested but %s (sub-group size %llu)\n", why, (unsigned long long)sg); exit(2); }
        }
    }
    if (buildProgram(dv, di, src, srcLen, 0, groupSize, o->extraOpts) != 0) { printf("FAIL: kernel.cl did not compile\n"); exit(2); }
    dv->exchange = 0;
    if (dv->subGroupSize) snprintf(dv->exchangeNote, sizeof(dv->exchangeNote), "%s (%s; queried sub-group size %llu)", exchangeName(0), why, (unsigned long long)dv->subGroupSize);
    else snprintf(dv->exchangeNote, sizeof(dv->exchangeNote), "%s (%s)", exchangeName(0), why);
}

// ---------------------------------------------------------------------------------------------
// Launch helpers

static size_t kernelMaxLocal(const Device* dv, cl_kernel k, const DeviceInfo* di, size_t want) {
    size_t wg = 0;
    (void)dv;
    if (clGetKernelWorkGroupInfo(k, di->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(wg), &wg, NULL) != CL_SUCCESS || wg == 0) wg = want;
    if (wg > di->maxWorkGroup) wg = di->maxWorkGroup;
    return wg < want ? wg : want;
}

/* Object accounting for the serve loop's stats line (4 October 2026, after the gfx1036 fault): every event and buffer
 * the serve path creates or releases is counted here, so a leak shows as a growing "live" count long before a runtime
 * limit is hit. The bench path does not count. */
static unsigned long gEvCreated = 0, gEvReleased = 0, gMemCreated = 0, gMemReleased = 0;
static void countRelease(cl_event ev) { clReleaseEvent(ev); ++gEvReleased; }

static cl_event launch1D(const Device* dv, cl_kernel k, size_t global, size_t local) {
    cl_event ev = NULL;
    size_t g = ((global + local - 1) / local) * local;
    CL_CHECK(clEnqueueNDRangeKernel(dv->q, k, 1, NULL, &g, &local, 0, NULL, &ev));
    ++gEvCreated;
    return ev;
}

static cl_event launchHash(const Device* dv, cl_mem ds, cl_mem out, cl_uint baseNonce, cl_uint mask, size_t nonces, size_t groupSize) {
    CL_CHECK(clSetKernelArg(dv->kHash, 0, sizeof(cl_mem), &ds));
    CL_CHECK(clSetKernelArg(dv->kHash, 1, sizeof(cl_mem), &out));
    CL_CHECK(clSetKernelArg(dv->kHash, 2, sizeof(cl_uint), &baseNonce));
    CL_CHECK(clSetKernelArg(dv->kHash, 3, sizeof(cl_uint), &mask));
    return launch1D(dv, dv->kHash, nonces, groupSize);
}

static void readWords(const Device* dv, cl_mem buf, size_t wordIndex, size_t nWords, uint32_t* dst) {
    CL_CHECK(clEnqueueReadBuffer(dv->q, buf, CL_TRUE, wordIndex * 4u, nWords * 4u, dst, 0, NULL, NULL));
}

static int compareWarp(const uint64_t* got, const uint64_t* want, uint32_t base, const char* how) {
    int bad = 0, first = -1, l;
    for (l = 0; l < 32; ++l) if (got[l] != want[l]) { if (first < 0) first = l; ++bad; }
    if (bad == 0) printf("verify warp base %u (nonces %u..%u) %s: PASS\n", base, base, base + 31u, how);
    else printf("verify warp base %u (nonces %u..%u) %s: FAIL  %d of 32 lanes differ, first lane %d: device=%016llx expected=%016llx\n",
                base, base, base + 31u, how, bad, first, (unsigned long long)got[first], (unsigned long long)want[first]);
    return bad == 0;
}

// ---------------------------------------------------------------------------------------------
// Cache (memory-hard packs)

#if IGNEUM_DATASET_MODE == 1
static cl_mem gCache = NULL;
static double gCacheFillFirstMs = 0, gCacheFillSecondMs = 0, gCacheHostMs = 0;
static int gCachePass = 0;

static int setupCache(Device* dv, const DeviceInfo* di) {
    size_t bytes = (size_t)CACHE_WORDS_HOST * 4u;
    cl_int err = 0;
    cl_uint nSeg = IGNEUM_CACHE_SEGMENTS;
    size_t local;
    int pass;
    uint32_t* dev;
    uint32_t seg;
    double h0;
    int same, fnvOk, headOk, lastOk;
    uint64_t fnv;
    gCache = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, bytes, NULL, &err);
    CL_CHECK_ERR(err, "clCreateBuffer cache");
    local = kernelMaxLocal(dv, dv->kCacheFill, di, 256);
    for (pass = 0; pass < 2; ++pass) {
        cl_event ev;
        CL_CHECK(clSetKernelArg(dv->kCacheFill, 0, sizeof(cl_mem), &gCache));
        CL_CHECK(clSetKernelArg(dv->kCacheFill, 1, sizeof(cl_uint), &nSeg));
        ev = launch1D(dv, dv->kCacheFill, nSeg, local);
        CL_CHECK(clFinish(dv->q));
        if (pass == 0) gCacheFillFirstMs = eventMs(ev); else gCacheFillSecondMs = eventMs(ev);
        countRelease(ev);
    }
    printf("cache fill (device): %.2f ms first, %.2f ms second (%u chains x %u ChaCha blocks, %u MiB, work-group %llu)\n",
           gCacheFillFirstMs, gCacheFillSecondMs, (unsigned)IGNEUM_CACHE_SEGMENTS, 1u << IGNEUM_CACHE_SEGMENT_LOG2_LINES,
           (unsigned)(bytes >> 20), (unsigned long long)local);

    hCache = (uint32_t*)malloc(bytes);
    h0 = wallMs();
    for (seg = 0; seg < IGNEUM_CACHE_SEGMENTS; ++seg) mh_cache_segment(hCache, seg);
    gCacheHostMs = wallMs() - h0;
    printf("cache fill (host, one thread): %.1f ms\n", gCacheHostMs);

    dev = (uint32_t*)malloc(bytes);
    readWords(dv, gCache, 0, CACHE_WORDS_HOST, dev);
    same = memcmp(dev, hCache, bytes) == 0;
    fnv = fnv1a64(hCache, bytes);
    fnvOk = (fnv == IGNEUM_CACHE_FNV64);
    headOk = memcmp(hCache, IGNEUM_CACHE_HEAD, 64) == 0;
    lastOk = memcmp(hCache + CACHE_WORDS_HOST - 16u, IGNEUM_CACHE_LAST, 64) == 0;
    if (!same) {
        uint32_t i;
        for (i = 0; i < CACHE_WORDS_HOST; ++i) if (dev[i] != hCache[i]) {
            printf("  cache[%u]: device 0x%08x host 0x%08x (first difference)\n", i, dev[i], hCache[i]); break;
        }
    }
    free(dev);
    gCachePass = same && fnvOk && headOk && lastOk;
    printf("cache check: %s  (device == host all %u words %s, host FNV-1a 64 %016llx vs Mac %016llx %s, head 16 vs Mac %s, last line vs Mac %s)\n",
           gCachePass ? "PASS" : "FAIL", CACHE_WORDS_HOST, same ? "PASS" : "FAIL",
           (unsigned long long)fnv, (unsigned long long)IGNEUM_CACHE_FNV64, fnvOk ? "PASS" : "FAIL",
           headOk ? "PASS" : "FAIL", lastOk ? "PASS" : "FAIL");
    return gCachePass;
}
#endif

// ---------------------------------------------------------------------------------------------
// One dataset size: fill or build, self-test, vectors, bench

typedef struct {
    int mib;
    uint32_t words;
    double fillFirstMs, fillSecondMs;
    int dsPass;
    int vecChecked, vecPass;
    double gpuMs, wallMsTimed;
    double hashesPerSec, gbps;
} SizeResult;

static SizeResult runSize(Device* dv, const DeviceInfo* di, const Options* o, int mib, cl_mem dOut, uint32_t nonces) {
    SizeResult r;
    uint64_t bytes = (uint64_t)mib << 20;
    uint32_t mask;
    int atPackSize;
    cl_int err = 0;
    cl_mem dDs;
    int pass;
    size_t groupSize = 32u * (size_t)o->groupWarps;
    uint64_t got[32];
    double w0, w1, t0, t1, total;
    cl_event* evs;
    int b;
    memset(&r, 0, sizeof(r));
    r.mib = mib;
    r.words = (uint32_t)(bytes / 4ull);
    mask = r.words - 1u;
    atPackSize = (r.words == (1u << IGNEUM_DATASET_LOG2));
    printf("\n=== dataset %d MiB (2^%d words, mask 0x%08x)%s ===\n", mib, log2u32(r.words), mask,
           atPackSize ? "" : "  [not the pack size: vectors skipped, dataset head and random points still checked]");
    if ((uint64_t)di->maxAlloc < bytes) {
        printf("FAIL: CL_DEVICE_MAX_MEM_ALLOC_SIZE is %llu MiB, the dataset needs %d MiB in one buffer\n", (unsigned long long)(di->maxAlloc >> 20), mib);
        exit(2);
    }
    dDs = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, (size_t)bytes, NULL, &err);
    CL_CHECK_ERR(err, "clCreateBuffer dataset");

    // Fill (or build) twice: the Mac showed a first-touch cost on the first fill of a process.
    for (pass = 0; pass < 2; ++pass) {
        cl_event ev;
#if IGNEUM_DATASET_MODE == 1
        cl_uint nItems = r.words / 16u;
        size_t local = kernelMaxLocal(dv, dv->kBuild, di, 256);
        CL_CHECK(clSetKernelArg(dv->kBuild, 0, sizeof(cl_mem), &dDs));
        CL_CHECK(clSetKernelArg(dv->kBuild, 1, sizeof(cl_mem), &gCache));
        CL_CHECK(clSetKernelArg(dv->kBuild, 2, sizeof(cl_uint), &nItems));
        ev = launch1D(dv, dv->kBuild, nItems, local);
#else
        cl_uint n = r.words, d0 = IGNEUM_DAY0, d1 = IGNEUM_DAY1;
        size_t local = kernelMaxLocal(dv, dv->kFill, di, 256);
        CL_CHECK(clSetKernelArg(dv->kFill, 0, sizeof(cl_mem), &dDs));
        CL_CHECK(clSetKernelArg(dv->kFill, 1, sizeof(cl_uint), &n));
        CL_CHECK(clSetKernelArg(dv->kFill, 2, sizeof(cl_uint), &d0));
        CL_CHECK(clSetKernelArg(dv->kFill, 3, sizeof(cl_uint), &d1));
        ev = launch1D(dv, dv->kFill, n, local);
#endif
        CL_CHECK(clFinish(dv->q));
        if (pass == 0) r.fillFirstMs = eventMs(ev); else r.fillSecondMs = eventMs(ev);
        countRelease(ev);
    }
#if IGNEUM_DATASET_MODE == 1
    printf("dataset build (memory-hard, from the cache): %.2f ms first, %.2f ms second -> %.1f M items/s, %.2f G cache-line reads/s (second, device time)\n",
           r.fillFirstMs, r.fillSecondMs, (double)(r.words / 16u) / 1e6 / (r.fillSecondMs / 1000.0),
           (double)(r.words / 16u) * (double)IGNEUM_ITEM_ROUNDS / 1e9 / (r.fillSecondMs / 1000.0));
#else
    printf("dataset fill: %.2f ms first, %.2f ms second -> %.0f GB/s write (second, device time)\n",
           r.fillFirstMs, r.fillSecondMs, (double)bytes / 1e9 / (r.fillSecondMs / 1000.0));
#endif

    // Dataset self-test: head 16 (any size), element [MASK] (pack size only), 64 pseudo-random points vs the host
    // formula (closed form) or the host derivation from the host cache (memory-hard), and the Mac's 64 sampled words.
    {
        uint32_t head[16];
        int badHead = 0, i, k, badRnd = 0, badSample = 0, nSample = 0, lastOk = 1;
        const char* lastText = "skipped";
        uint64_t s = 0x9E3779B97F4A7C15ull ^ (uint64_t)r.words;
        readWords(dv, dDs, 0, 16, head);
        for (i = 0; i < 16; ++i) if (head[i] != IGNEUM_DS_HEAD[i]) {
            if (badHead == 0) printf("  dataset[%d] = 0x%08x, expected 0x%08x\n", i, head[i], IGNEUM_DS_HEAD[i]);
            ++badHead;
        }
        if (atPackSize) {
            uint32_t last = 0;
            readWords(dv, dDs, IGNEUM_DS_LAST_INDEX, 1, &last);
            lastOk = (last == IGNEUM_DS_LAST);
            lastText = lastOk ? "PASS" : "FAIL";
            if (!lastOk) printf("  dataset[%u] = 0x%08x, expected 0x%08x\n", IGNEUM_DS_LAST_INDEX, last, IGNEUM_DS_LAST);
        }
        for (k = 0; k < 64; ++k) {
            uint64_t z;
            uint32_t idx, v = 0, want;
            s += 0x9E3779B97F4A7C15ull;
            z = s;
            z = (z ^ (z >> 30)) * 0xBF58476D1CE4E5B9ull;
            z = (z ^ (z >> 27)) * 0x94D049BB133111EBull;
            z ^= z >> 31;
            idx = (uint32_t)z & mask;
            readWords(dv, dDs, idx, 1, &v);
#if IGNEUM_DATASET_MODE == 1
            want = host_ds_word(idx);
#else
            want = host_ds_elem(idx, IGNEUM_DAY0, IGNEUM_DAY1);
#endif
            if (v != want) {
                if (badRnd == 0) printf("  dataset[%u] = 0x%08x, host %s 0x%08x\n", idx, v, IGNEUM_DATASET_MODE == 1 ? "derivation" : "formula", want);
                ++badRnd;
            }
        }
#ifdef IGNEUM_DS_SAMPLES
        for (k = 0; k < IGNEUM_DS_SAMPLES; ++k) {
            uint32_t v = 0;
            if (IGNEUM_DS_SAMPLE_INDEX[k] > mask) continue;
            ++nSample;
            readWords(dv, dDs, IGNEUM_DS_SAMPLE_INDEX[k], 1, &v);
            if (v != IGNEUM_DS_SAMPLE_VALUE[k]) {
                if (badSample == 0) printf("  dataset[%u] = 0x%08x, Mac 0x%08x\n", IGNEUM_DS_SAMPLE_INDEX[k], v, IGNEUM_DS_SAMPLE_VALUE[k]);
                ++badSample;
            }
        }
#endif
        r.dsPass = (badHead == 0 && lastOk && badRnd == 0 && badSample == 0);
        printf("dataset self-test: %s  (head 16 vs Mac %s, element [MASK] vs Mac %s, 64 random points vs host %s %s, %d Mac samples %s)\n",
               r.dsPass ? "PASS" : "FAIL", badHead == 0 ? "PASS" : "FAIL", lastText,
               IGNEUM_DATASET_MODE == 1 ? "derivation" : "formula", badRnd == 0 ? "PASS" : "FAIL",
               nSample, nSample == 0 ? "none in pack" : (badSample == 0 ? "PASS" : "FAIL"));
    }

    // Vectors, standalone: one 32-item work-group per base nonce, exactly like the Mac cross-check.
    if (atPackSize) {
        int w;
        r.vecChecked = 1; r.vecPass = 1;
        for (w = 0; w < IGNEUM_VEC_WARPS; ++w) {
            cl_event ev;
            if (o->groupWarps != 1) {
                // reqd_work_group_size pins the hash kernel to 32 x group-warps items; run one full group and read its first 32.
                ev = launchHash(dv, dDs, dOut, IGNEUM_VEC_BASE[w], mask, groupSize, groupSize);
            } else {
                ev = launchHash(dv, dDs, dOut, IGNEUM_VEC_BASE[w], mask, 32u, 32u);
            }
            CL_CHECK(clFinish(dv->q));
            countRelease(ev);
            CL_CHECK(clEnqueueReadBuffer(dv->q, dOut, CL_TRUE, 0, sizeof(got), got, 0, NULL, NULL));
            r.vecPass = compareWarp(got, IGNEUM_VEC_OUT[w], IGNEUM_VEC_BASE[w], o->groupWarps == 1 ? "standalone, 1 unit/work-group" : "standalone, first unit of one work-group") && r.vecPass;
        }
    } else {
        printf("vectors: skipped (the pack's vectors are for %d MiB)\n", packMib());
    }

    // Warm-up batch at base nonce 0. With the default 2^24 nonces it contains all three vector warps,
    // so the bench configuration itself (work-group = 32 x group-warps) is also checked bit for bit.
    w0 = wallMs();
    {
        cl_event ev = launchHash(dv, dDs, dOut, 0u, mask, nonces, groupSize);
        CL_CHECK(clFinish(dv->q));
        countRelease(ev);
    }
    w1 = wallMs();
    printf("warm-up batch: %u hashes in %.2f ms wall\n", nonces, w1 - w0);
    if (atPackSize) {
        // Fingerprint of every output in the batch, so two runs (two devices, two exchange paths, the CPU emulator at the
        // same --batch-log2) can be compared for all nonces, not only the vector warps.
        uint64_t* all = (uint64_t*)malloc((size_t)nonces * sizeof(uint64_t));
        uint64_t fp = 0xcbf29ce484222325ull;
        size_t i;
        CL_CHECK(clEnqueueReadBuffer(dv->q, dOut, CL_TRUE, 0, (size_t)nonces * sizeof(uint64_t), all, 0, NULL, NULL));
        for (i = 0; i < (size_t)nonces * 8u; ++i) { fp ^= ((const uint8_t*)all)[i]; fp *= 0x100000001b3ull; }
        free(all);
        printf("batch fingerprint (FNV-1a 64 of 2^%d outputs at base nonce 0): %016llx\n", o->batchLog2, (unsigned long long)fp);
    }
    if (atPackSize) {
        char how[64];
        int w;
        snprintf(how, sizeof(how), "in batch, %d unit(s)/work-group", o->groupWarps);
        for (w = 0; w < IGNEUM_VEC_WARPS; ++w) {
            if ((uint64_t)IGNEUM_VEC_BASE[w] + 32ull > (uint64_t)nonces) {
                printf("verify warp base %u in batch: skipped (batch has %u nonces)\n", IGNEUM_VEC_BASE[w], nonces);
                continue;
            }
            CL_CHECK(clEnqueueReadBuffer(dv->q, dOut, CL_TRUE, (size_t)IGNEUM_VEC_BASE[w] * 8u, sizeof(got), got, 0, NULL, NULL));
            r.vecPass = compareWarp(got, IGNEUM_VEC_OUT[w], IGNEUM_VEC_BASE[w], how) && r.vecPass;
        }
    }

    // Timed batches. Base nonces (b * nonces) mod 2^32, as in the Metal and CUDA runs. Device time is the span from
    // the first batch's start to the last batch's end from event profiling, like cudaEvent elapsed time.
    evs = (cl_event*)calloc((size_t)o->batches, sizeof(cl_event));
    t0 = wallMs();
    for (b = 1; b <= o->batches; ++b) {
        uint32_t base = (uint32_t)((uint64_t)b * (uint64_t)nonces);
        evs[b - 1] = launchHash(dv, dDs, dOut, base, mask, nonces, groupSize);
    }
    CL_CHECK(clFinish(dv->q));
    t1 = wallMs();
    total = (double)nonces * (double)o->batches;
    r.gpuMs = spanMs(evs[0], evs[o->batches - 1]);
    for (b = 0; b < o->batches; ++b) clReleaseEvent(evs[b]);
    free(evs);
    r.wallMsTimed = t1 - t0;
    if (r.gpuMs < 0 && !o->timeWall) { printf("NOTE: event profiling unavailable on this runtime; the rate uses wall time\n"); ((Options*)o)->timeWall = 1; }
    r.hashesPerSec = total / ((o->timeWall ? r.wallMsTimed : r.gpuMs) / 1000.0);
    r.gbps = r.hashesPerSec * (double)IGNEUM_LOADS_PER_HASH * 4.0 / 1e9;
    printf("timed: %d batches x %u hashes = %.0f hashes (rate below from %s time)\n", o->batches, nonces, total, o->timeWall ? "wall" : "device event");
    printf("  device %.2f ms -> %.3f Mhash/s%s\n", r.gpuMs, total / (r.gpuMs / 1000.0) / 1e6, o->timeWall ? "  (event profiling, not used for the rate on this platform)" : "");
    printf("  wall   %.2f ms -> %.3f Mhash/s\n", r.wallMsTimed, total / (r.wallMsTimed / 1000.0) / 1e6);
    printf("  rate   %.3f Mhash/s (%.0f hashes/s), %.2f GB/s useful (loads x 4 B)\n", r.hashesPerSec / 1e6, r.hashesPerSec, r.gbps);

    CL_CHECK(clReleaseMemObject(dDs));
    return r;
}

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


/* ---------------------------------------------------------------------------------------------
 * Serve mode: GPU worker for igneum-miner --worker (3 October 2026)
 *
 * Protocol, one line each. Only "found", "done" and "error" are parsed by the miner; every other line is logged.
 *   stdin:  job <job_id> <header_prehash_hex 64> <target_hex 16> <nonce_start u64> <nonce_count u64> <epoch_seed_hex 64> <day_seed_hex>
 *           quit
 *   stdout: ready opencl <device>
 *           found <job_id> <nonce u64> <hash_hex 16>      every nonce whose 64-bit hash is <= target
 *           done <job_id> <hashes> <ms>                     end of the job (wall ms)
 *           error <job_id> <text>
 *           prepare <epoch_seed_hex 64> <day_seed_hex> <pack_dir>   build <pack_dir>/kernel_bound.cl (the pack the miner wrote
 *                                                                  for those seeds) plus its cache and dataset in the background
 *   stdout: ready opencl <device> ... prepare 1
 *           prepared <epoch_seed_hex> <day_seed_hex> <ms> ...   the pair is resident; a job on it switches instantly
 *           prepare-failed <epoch_seed_hex> <day_seed_hex> <text>
 * The pack's program.h is compiled in and kernel_bound.cl is built at runtime, so at start this worker serves exactly
 * one epoch seed and one day seed: the pack's. The next pair arrives through `prepare`: the miner writes the pack for the
 * prepared seeds (igneum-miner --prepare-packs <dir>) and names its directory; a background thread builds that pack's
 * kernel_bound.cl with the same build options, fills its cache and builds its dataset on a second queue while jobs on
 * the current pair keep running (at most two pairs resident: the current one and the prepared one; the old pair is
 * released after the first job on the new one). A job for seeds that are neither the current nor the prepared pair is
 * answered with an error naming both. Without prepare, re-export the pack with `igneum-miner export-pack <node> <dir>`
 * and rebuild. A prepared pair's cache is not cross-checked against a host fill (memhard.h is compiled in for the
 * original day); the miner's CPU re-check of every found nonce covers it. The init words of a dispatch are
 * seed_words_from_bytes("igneum-block/" || prehash || nonce_hi_le32), written to a small buffer that is the fifth
 * argument of igneum_hash_bound; the lane nonce is baseNonce + gid as in the bench kernel. The exchange rule of
 * WAVEFRONT.md applies unchanged (the bound kernel has the same body and the same IGNEUM_EXCHANGE build).
 */

static void seedWordsFromBytes(const uint8_t* b, size_t n, uint32_t out[8]) {
    uint64_t salt;
    for (salt = 0; salt < 4; ++salt) {
        uint64_t h = 0xcbf29ce484222325ull ^ (salt * 0x9E3779B97F4A7C15ull);
        size_t i;
        for (i = 0; i < n; ++i) { h ^= b[i]; h *= 0x100000001b3ull; }
        h ^= h >> 33; h *= 0xff51afd7ed558ccdull; h ^= h >> 33;
        out[2 * salt] = (uint32_t)h;
        out[2 * salt + 1] = (uint32_t)(h >> 32);
    }
}

static int unhexBuf(const char* s, uint8_t* out, size_t cap, size_t* len) {
    size_t n = strlen(s), i;
    if (n % 2 || n / 2 > cap) return 0;
    for (i = 0; i < n; i += 2) {
        unsigned v = 0;
        char two[3] = { s[i], s[i + 1], 0 };
        if (sscanf(two, "%2x", &v) != 1) return 0;
        out[i / 2] = (uint8_t)v;
    }
    *len = n / 2;
    return 1;
}

/* Generic serve mode (--pack, 4 October 2026): the sizes and seeds come from the pack directory read at run time, not
 * from the compiled-in program.h, so one prebuilt exe serves every pack. The compiled-in values are the defaults. */
static PfPack gPack;
static int gGeneric = 0;
/* Variant 5 of the read-width experiment (5 October 2026): the scratch arena of a persistent-warp pack, its warp count
 * and the running tag salt; set by --bench-pack before the self-test. Serve mode does not support these packs. */
static cl_mem gScratch = NULL;
static cl_uint gScratchWarps = 0;
static cl_uint gSalt = 1;
/* Sets the three extra arguments of a variant-5 kernel (after the five of igneum_hash_bound) for `units` units. */
static cl_int setScratchArgs(cl_kernel k, cl_uint firstArg, cl_uint units) {
    cl_int e = clSetKernelArg(k, firstArg, sizeof(cl_mem), &gScratch);
    if (e == CL_SUCCESS) e = clSetKernelArg(k, firstArg + 1, sizeof(cl_uint), &units);
    if (e == CL_SUCCESS) e = clSetKernelArg(k, firstArg + 2, sizeof(cl_uint), &gSalt);
    gSalt += units;
    return e;
}
static uint32_t gServeWords = 1u << IGNEUM_DATASET_LOG2;
#if IGNEUM_DATASET_MODE == 1
static uint32_t gServeCacheWords = 1u << IGNEUM_CACHE_LOG2_WORDS;
static uint32_t gServeSegments = IGNEUM_CACHE_SEGMENTS;
#endif

#if IGNEUM_DATASET_MODE == 1
/* One resident (program, cache, dataset) triple for a seed pair. The first one is the compiled-in pack on the main
 * queue (or, with --pack, the directory's pack); prepared ones are built from a pack directory on their own queue. */
typedef struct {
    char epochHex[65];
    char dayHex[512];
    uint32_t sw[8], kw[8];      /* seed words and key words, for job matching */
    cl_program prog;
    cl_kernel kHashBound, kCacheFill, kBuild;
    cl_mem cache, ds;
    /* hot-table experiment (5 October 2026, docs/plans/hot-table.md): the epoch's table, filled on the device by the
     * pack's igneum_hot_fill, the argument after the init words; hotWords 0 for a pack without one */
    cl_kernel kHotFill;
    cl_mem hot;
    uint32_t hotWords, hotSegments, hotMb, hotSlots;
    double hotMs;
    double buildMs, cacheMs, datasetMs, checkMs;
    char check[1024];           /* the self-test verdict (packfile.h), one line */
    int checked;
    char programClass[8];       /* Counter ASIC 2.0: the pack's class ("v2" or "v3") and era seed, from packfile.h */
    char eraHex[65];
} ServePair;

static int hexEq(const char* a, const char* b) {
    size_t i;
    if (strlen(a) != strlen(b)) return 0;
    for (i = 0; a[i]; ++i) if (tolower((unsigned char)a[i]) != tolower((unsigned char)b[i])) return 0;
    return 1;
}

/* 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. The compiled-in placeholder pack
 * (no --pack, no prepared pair) has no seed hex; it keeps the word comparison. */
static int pairIs(const ServePair* p, const char* epochHex, const char* dayHex, const uint32_t sw[8], const uint32_t kw[8]) {
    if (!p) return 0;
    if (p->epochHex[0]) return hexEq(p->epochHex, epochHex) && hexEq(p->dayHex, dayHex);
    return memcmp(sw, p->sw, 32) == 0 && memcmp(kw, p->kw, 32) == 0;
}

/* Counter ASIC 2.0: a job that names a class (and an era) belongs to a pair of that class (and era) only. */
static int pairIsClass(const ServePair* p, const char* epochHex, const char* dayHex, const uint32_t sw[8], const uint32_t kw[8], const char* cls, const char* era) {
    char why[256];
    return pairIs(p, epochHex, dayHex, sw, kw) && pf_pack_class_ok(p->programClass[0] ? p->programClass : "v2", p->eraHex, cls, era, 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(char** f, int* nf, char* cls, size_t clsCap, char* era, size_t eraCap) {
    cls[0] = 0; era[0] = 0;
    while (*nf > 0 && pf_class_token(f[*nf - 1], cls, clsCap, era, eraCap)) --*nf;
}

static void releasePair(ServePair* p) {
    if (!p) return;
    if (p->ds) { clReleaseMemObject(p->ds); ++gMemReleased; }
    if (p->cache) { clReleaseMemObject(p->cache); ++gMemReleased; }
    if (p->hot) { clReleaseMemObject(p->hot); ++gMemReleased; }
    if (p->kHotFill) clReleaseKernel(p->kHotFill);
    if (p->kHashBound) clReleaseKernel(p->kHashBound);
    if (p->kCacheFill) clReleaseKernel(p->kCacheFill);
    if (p->kBuild) clReleaseKernel(p->kBuild);
    if (p->prog) clReleaseProgram(p->prog);
    free(p);
}

/* The prepare request and its result, handed between the main loop and the prepare thread. */
typedef struct {
    Device* dv;
    const DeviceInfo* di;
    char epochHex[65];
    char dayHex[512];
    char packDir[1024];
    char wantClass[8];          /* the class and era the prepare line named (empty: any) */
    char wantEra[65];
    uint32_t words, cacheWords, segments;
    char error[512];
    ServePair* result;          /* set by the thread on success */
    volatile int done;          /* 1 when the thread has finished (success or failure) */
    double t0, doneAt;
} PrepareTask;

static void prepareFail(PrepareTask* t, const char* what, cl_int err) {
    snprintf(t->error, sizeof(t->error), "%s (%s)", what, clErrName(err));
}

/* The self-test of packfile.h on a pair: cache head, last line and FNV-1a 64 over the whole cache, dataset head, last
 * word and samples, and the vector warps through igneum_hash_bound with the pack's own seed words as init words (that
 * is igneum_hash of kernel.cl). The verdict goes to p->check. A pack directory without a readable vectors.h is not
 * tested (p->checked stays 0: the miner's CPU re-check still covers every found nonce); a FAIL is an error. */
static int pairSelfTest(Device* dv, const DeviceInfo* di, cl_command_queue q, ServePair* p, uint32_t words, uint32_t cacheWords, const char* packDir, char* err, size_t errCap) {
    PfPack pk;
    char perr[256];
    uint32_t cacheHead[16], cacheLast[16], dsHead[16], dsLast = 0;
    uint32_t hotHead[16], hotLast[16];
    uint64_t hotFnv = 0;
    uint32_t samples[PF_MAX_SAMPLES];
    uint64_t vec[PF_MAX_WARPS * 32];
    uint64_t fnv;
    uint32_t* whole;
    cl_mem out = NULL, init = NULL;
    cl_int e = 0;
    int w, i, ok;
    cl_uint extra = 5;   /* the first argument after the five of igneum_hash_bound: the hot table, then the scratch triple */
    size_t local = (size_t)dv->groupSize, g = (size_t)dv->groupSize;   /* one work-group; its first 32 lanes are the warp */
    double tb = wallMs();
    (void)di;
    p->checked = 0;
    perr[0] = 0;
    if (!packDir || !packDir[0] || !pf_load(packDir, &pk, perr, sizeof(perr)) || !pk.haveVectors) {
        snprintf(p->check, sizeof(p->check), "self-test skipped (%s); the miner's CPU re-check covers every found nonce", (packDir && packDir[0]) ? (perr[0] ? perr : "no vectors.h in the pack") : "no pack directory");
        return 1;
    }
    whole = (uint32_t*)malloc((size_t)cacheWords * 4u);
    if (!whole) { snprintf(err, errCap, "self-test: no host memory for the cache read-back"); return 0; }
    e = clEnqueueReadBuffer(q, p->cache, CL_TRUE, 0, (size_t)cacheWords * 4u, whole, 0, NULL, NULL);
    if (e != CL_SUCCESS) { free(whole); snprintf(err, errCap, "self-test: cache read-back (%s)", clErrName(e)); return 0; }
    memcpy(cacheHead, whole, 64);
    memcpy(cacheLast, whole + cacheWords - 16u, 64);
    fnv = pf_fnv1a64(whole, (size_t)cacheWords * 4u);
    free(whole);
    e = clEnqueueReadBuffer(q, p->ds, CL_TRUE, 0, 64, dsHead, 0, NULL, NULL);
    if (e == CL_SUCCESS && pk.dsLastIndex < words) e = clEnqueueReadBuffer(q, p->ds, CL_TRUE, (size_t)pk.dsLastIndex * 4u, 4, &dsLast, 0, NULL, NULL);
    for (i = 0; i < pk.nSamples && e == CL_SUCCESS; ++i) { samples[i] = 0; if (pk.sampleIdx[i] < words) e = clEnqueueReadBuffer(q, p->ds, CL_TRUE, (size_t)pk.sampleIdx[i] * 4u, 4, &samples[i], 0, NULL, NULL); }
    if (e != CL_SUCCESS) { snprintf(err, errCap, "self-test: dataset read-back (%s)", clErrName(e)); return 0; }
    if (p->hot) {
        uint32_t* hw = (uint32_t*)malloc((size_t)p->hotWords * 4u);
        if (!hw) { snprintf(err, errCap, "self-test: no host memory for the hot table read-back"); return 0; }
        e = clEnqueueReadBuffer(q, p->hot, CL_TRUE, 0, (size_t)p->hotWords * 4u, hw, 0, NULL, NULL);
        if (e != CL_SUCCESS) { free(hw); snprintf(err, errCap, "self-test: hot table read-back (%s)", clErrName(e)); return 0; }
        memcpy(hotHead, hw, 64);
        memcpy(hotLast, hw + p->hotWords - 16u, 64);
        hotFnv = pf_fnv1a64(hw, (size_t)p->hotWords * 4u);
        free(hw);
    }
    out = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, g * sizeof(uint64_t), NULL, &e);
    if (e == CL_SUCCESS) { ++gMemCreated; init = clCreateBuffer(dv->ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 32, pk.seedw, &e); }
    if (e == CL_SUCCESS) ++gMemCreated;
    if (pk.persistent) {
        if (!gScratch) { snprintf(err, errCap, "self-test: a variant-5 pack (persistent warps) needs --bench-pack (serve mode does not carry a scratch)"); return 0; }
        g = 32; local = 32;   /* one persistent warp runs the one unit */
    }
    for (w = 0; w < pk.vecWarps && e == CL_SUCCESS; ++w) {
        cl_uint base = pk.vecBase[w], mask = words - 1u;
        e = clSetKernelArg(p->kHashBound, 0, sizeof(cl_mem), &p->ds);
        if (e == CL_SUCCESS) e = clSetKernelArg(p->kHashBound, 1, sizeof(cl_mem), &out);
        if (e == CL_SUCCESS) e = clSetKernelArg(p->kHashBound, 2, sizeof(cl_uint), &base);
        if (e == CL_SUCCESS) e = clSetKernelArg(p->kHashBound, 3, sizeof(cl_uint), &mask);
        if (e == CL_SUCCESS) e = clSetKernelArg(p->kHashBound, 4, sizeof(cl_mem), &init);
        if (e == CL_SUCCESS && p->hot) { e = clSetKernelArg(p->kHashBound, 5, sizeof(cl_mem), &p->hot); extra = 6; }
        if (e == CL_SUCCESS && pk.persistent) e = setScratchArgs(p->kHashBound, extra, 1u);
        if (e == CL_SUCCESS) e = clEnqueueNDRangeKernel(q, p->kHashBound, 1, NULL, &g, &local, 0, NULL, NULL);
        if (e == CL_SUCCESS) e = clEnqueueReadBuffer(q, out, CL_TRUE, 0, 32 * sizeof(uint64_t), &vec[w * 32], 0, NULL, NULL);
    }
    if (out) { clReleaseMemObject(out); ++gMemReleased; }
    if (init) { clReleaseMemObject(init); ++gMemReleased; }
    if (e != CL_SUCCESS) { snprintf(err, errCap, "self-test: vector warp (%s)", clErrName(e)); return 0; }
    ok = pf_selftest(&pk, cacheHead, cacheLast, fnv, dsHead, dsLast, samples, vec, p->hot ? hotHead : NULL, p->hot ? hotLast : NULL, hotFnv, p->check, sizeof(p->check));
    p->checked = 1;
    p->checkMs = wallMs() - tb;
    if (!ok) { snprintf(err, errCap, "%s", p->check); return 0; }
    return 1;
}

/* Fills the pair's cache and builds its dataset on queue q (the pair's kernels must exist), then self-tests it.
 * Shared by the first pair of --pack and by every prepared pair. Returns 1, or 0 with err set. */
static int pairBuffers(Device* dv, const DeviceInfo* di, cl_command_queue q, ServePair* p, uint32_t words, uint32_t cacheWords, uint32_t segments, const char* packDir, char* err, size_t errCap) {
    cl_int e = 0;
    cl_uint nSeg = segments, nItems = words / 16u;
    size_t local, g, bytes = (size_t)cacheWords * 4u;
    double tb = wallMs();
    p->cache = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, bytes, NULL, &e);
    if (e != CL_SUCCESS) { snprintf(err, errCap, "clCreateBuffer cache (%s)", clErrName(e)); return 0; }
    ++gMemCreated;
    local = kernelMaxLocal(dv, p->kCacheFill, di, 256);
    g = ((nSeg + local - 1) / local) * local;
    e = clSetKernelArg(p->kCacheFill, 0, sizeof(cl_mem), &p->cache);
    if (e == CL_SUCCESS) e = clSetKernelArg(p->kCacheFill, 1, sizeof(cl_uint), &nSeg);
    if (e == CL_SUCCESS) e = clEnqueueNDRangeKernel(q, p->kCacheFill, 1, NULL, &g, &local, 0, NULL, NULL);
    if (e == CL_SUCCESS) e = clFinish(q);
    if (e != CL_SUCCESS) { snprintf(err, errCap, "cache fill (%s)", clErrName(e)); return 0; }
    p->cacheMs = wallMs() - tb;
    tb = wallMs();
    p->ds = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, (size_t)words * 4u, NULL, &e);
    if (e != CL_SUCCESS) { snprintf(err, errCap, "clCreateBuffer dataset (%s)", clErrName(e)); return 0; }
    ++gMemCreated;
    local = kernelMaxLocal(dv, p->kBuild, di, 256);
    g = ((nItems + local - 1) / local) * local;
    e = clSetKernelArg(p->kBuild, 0, sizeof(cl_mem), &p->ds);
    if (e == CL_SUCCESS) e = clSetKernelArg(p->kBuild, 1, sizeof(cl_mem), &p->cache);
    if (e == CL_SUCCESS) e = clSetKernelArg(p->kBuild, 2, sizeof(cl_uint), &nItems);
    if (e == CL_SUCCESS) e = clEnqueueNDRangeKernel(q, p->kBuild, 1, NULL, &g, &local, 0, NULL, NULL);
    if (e == CL_SUCCESS) e = clFinish(q);
    if (e != CL_SUCCESS) { snprintf(err, errCap, "dataset build (%s)", clErrName(e)); return 0; }
    p->datasetMs = wallMs() - tb;
    if (p->kHotFill && p->hotWords) {
        /* hot-table experiment: the epoch's table from the pack's own fill kernel (one work-item per segment) */
        cl_uint nSegH = p->hotSegments;
        tb = wallMs();
        p->hot = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, (size_t)p->hotWords * 4u, NULL, &e);
        if (e != CL_SUCCESS) { snprintf(err, errCap, "clCreateBuffer hot table (%s)", clErrName(e)); return 0; }
        ++gMemCreated;
        local = kernelMaxLocal(dv, p->kHotFill, di, 256);
        g = ((nSegH + local - 1) / local) * local;
        e = clSetKernelArg(p->kHotFill, 0, sizeof(cl_mem), &p->hot);
        if (e == CL_SUCCESS) e = clSetKernelArg(p->kHotFill, 1, sizeof(cl_uint), &nSegH);
        if (e == CL_SUCCESS) e = clEnqueueNDRangeKernel(q, p->kHotFill, 1, NULL, &g, &local, 0, NULL, NULL);
        if (e == CL_SUCCESS) e = clFinish(q);
        if (e != CL_SUCCESS) { snprintf(err, errCap, "hot table fill (%s)", clErrName(e)); return 0; }
        p->hotMs = wallMs() - tb;
    }
    return pairSelfTest(dv, di, q, p, words, cacheWords, packDir, err, errCap);
}

/* The hot-table kernel and sizes of a pair whose program came from a pack directory (none when the pack has no hot table). */
static int pairHotKernel(ServePair* p, const char* packDir, char* err, size_t errCap) {
    PfPack pk;
    char perr[256];
    cl_int e = 0;
    if (!packDir || !packDir[0] || !pf_load(packDir, &pk, perr, sizeof(perr)) || !pk.hotMb) return 1;
    p->kHotFill = clCreateKernel(p->prog, "igneum_hot_fill", &e);
    if (e != CL_SUCCESS) { snprintf(err, errCap, "clCreateKernel igneum_hot_fill (%s): the pack says IGNEUM_HOT_MB %u but its kernel source has no hot fill", clErrName(e), pk.hotMb); return 0; }
    p->hotWords = pk.hotWords; p->hotSegments = pk.hotSegments; p->hotMb = pk.hotMb; p->hotSlots = pk.hotSlots;
    return 1;
}

/* Builds the pair for a prepare request. Runs on its own thread with its own command queue. */
static void prepareRun(PrepareTask* t) {
    cl_int err = 0;
    size_t srcLen = 0;
    char path[1200];
    char* src;
    ServePair* p = (ServePair*)calloc(1, sizeof(ServePair));
    cl_command_queue q = NULL;
    double tb;
    strncpy(p->epochHex, t->epochHex, 64); p->epochHex[64] = 0;
    strncpy(p->dayHex, t->dayHex, sizeof(p->dayHex) - 1);
    {
        /* Counter ASIC 2.0: the pack's class and era are read before anything is built; a pack of another class
         * than the prepare line names is refused here, so the miner exports one of the right class */
        PfPack pk; char perr[512]; char why[256];
        if (!pf_load(t->packDir, &pk, perr, sizeof(perr))) { snprintf(t->error, sizeof(t->error), "pack %s: %s", t->packDir, perr); releasePair(p); t->doneAt = wallMs(); t->done = 1; return; }
        strncpy(p->programClass, pk.programClass, sizeof(p->programClass) - 1); strncpy(p->eraHex, pk.eraHex, sizeof(p->eraHex) - 1);
        if (!pf_pack_class_ok(p->programClass, p->eraHex, t->wantClass, t->wantEra, why, sizeof(why))) { snprintf(t->error, sizeof(t->error), "pack %s: %s", t->packDir, why); releasePair(p); t->doneAt = wallMs(); t->done = 1; return; }
    }
    snprintf(path, sizeof(path), "%s/kernel_bound.cl", t->packDir);
    src = readFile(path, &srcLen);
    if (!src) { snprintf(t->error, sizeof(t->error), "cannot read %s", path); releasePair(p); t->doneAt = wallMs(); t->done = 1; return; }
    tb = wallMs();
    p->prog = clCreateProgramWithSource(t->dv->ctx, 1, (const char**)&src, &srcLen, &err);
    free(src);
    if (err != CL_SUCCESS) { prepareFail(t, "clCreateProgramWithSource", err); releasePair(p); t->doneAt = wallMs(); t->done = 1; return; }
    err = clBuildProgram(p->prog, 1, &t->di->device, t->dv->buildOptions, NULL, NULL);
    if (err != CL_SUCCESS) {
        size_t logLen = 0;
        char* log;
        clGetProgramBuildInfo(p->prog, t->di->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &logLen);
        log = (char*)calloc(logLen + 1, 1);
        if (logLen) clGetProgramBuildInfo(p->prog, t->di->device, CL_PROGRAM_BUILD_LOG, logLen, log, NULL);
        snprintf(t->error, sizeof(t->error), "clBuildProgram failed (%s): %.300s", clErrName(err), log);
        free(log); releasePair(p); t->done = 1; return;
    }
    p->buildMs = wallMs() - tb;
    p->kHashBound = clCreateKernel(p->prog, "igneum_hash_bound", &err);
    if (err != CL_SUCCESS) { prepareFail(t, "clCreateKernel igneum_hash_bound (is this a kernel_bound.cl?)", err); releasePair(p); t->doneAt = wallMs(); t->done = 1; return; }
    p->kCacheFill = clCreateKernel(p->prog, "igneum_cache_fill", &err);
    if (err != CL_SUCCESS) { prepareFail(t, "clCreateKernel igneum_cache_fill", err); releasePair(p); t->doneAt = wallMs(); t->done = 1; return; }
    p->kBuild = clCreateKernel(p->prog, "igneum_build", &err);
    if (err != CL_SUCCESS) { prepareFail(t, "clCreateKernel igneum_build", err); releasePair(p); t->doneAt = wallMs(); t->done = 1; return; }
    if (!pairHotKernel(p, t->packDir, t->error, sizeof(t->error))) { releasePair(p); t->doneAt = wallMs(); t->done = 1; return; }
    q = clCreateCommandQueue(t->dv->ctx, t->di->device, 0, &err);
    if (err != CL_SUCCESS) { prepareFail(t, "clCreateCommandQueue (prepare)", err); releasePair(p); t->doneAt = wallMs(); t->done = 1; return; }
    if (!pairBuffers(t->dv, t->di, q, p, t->words, t->cacheWords, t->segments, t->packDir, t->error, sizeof(t->error))) { clReleaseCommandQueue(q); releasePair(p); t->doneAt = wallMs(); t->done = 1; return; }
    clReleaseCommandQueue(q);
    t->result = p;
    t->doneAt = wallMs();
    t->done = 1;
}

#ifdef _WIN32
static DWORD WINAPI prepareThreadMain(LPVOID arg) { prepareRun((PrepareTask*)arg); return 0; }
static int startPrepareThread(PrepareTask* t) { HANDLE h = CreateThread(NULL, 0, prepareThreadMain, t, 0, NULL); if (!h) return 0; CloseHandle(h); return 1; }
#else
static void* prepareThreadMain(void* arg) { prepareRun((PrepareTask*)arg); return NULL; }
static int startPrepareThread(PrepareTask* t) { pthread_t th; if (pthread_create(&th, NULL, prepareThreadMain, t) != 0) return 0; pthread_detach(th); return 1; }
#endif
#endif

/* --bench-pack (read-width experiment, 5 October 2026): the pack in --pack is built and self-tested exactly as the
 * first pair of --serve (pairBuffers: cache, dataset, cache FNV, dataset words, the vector warps through
 * igneum_hash_bound with the pack's seed words), then the bound kernel is timed over --batches dispatches of
 * 2^--batch-log2 nonces with device event time, and the 2^B outputs at base nonce 0 are fingerprinted (FNV-1a 64) so
 * the same pack can be compared bit for bit across vendors. A variant-5 pack is launched as --warps persistent warps
 * with a 1 MiB scratch each. One line per run starts with RESULT. */
static int runBenchPack(Device* dv, const DeviceInfo* di, const Options* o) {
#if IGNEUM_DATASET_MODE != 1
    (void)dv; (void)di; (void)o;
    printf("FAIL: --bench-pack needs a memory-hard placeholder pack\n");
    return 2;
#else
    const uint32_t words = gServeWords, mask = words - 1u;
    uint32_t nonces = 1u << o->batchLog2;
    size_t groupSize = 32 * (size_t)o->groupWarps, g;
    cl_int err = 0;
    cl_mem dOut, dInit;
    uint64_t* hOut;
    ServePair* cur;
    char perr[512], devName[256];
    double t0 = wallMs(), sum = 0, warmMs = 0;
    uint64_t fp = 0;
    int b, k;
    cl_uint warps = (cl_uint)(o->warps > 0 ? o->warps : 2048), units = nonces / 32u;
    if (!dv->kHashBound) { printf("FAIL: the kernel source has no igneum_hash_bound\n"); return 2; }
    if (gPack.persistent) {
        size_t arena;
        if (o->groupWarps != 1) { printf("FAIL: a variant-5 pack needs --group-warps 1 (one warp per work-group: the loop trip count must be uniform)\n"); return 2; }
        while (warps > 1 && units % warps != 0) warps >>= 1;
        arena = (size_t)warps * 32u * (size_t)gPack.scratchWordsPerLane * 4u;
        if ((uint64_t)arena > di->maxAlloc) { printf("FAIL: scratch arena %llu MiB exceeds the device's max alloc %llu MiB; lower --warps\n", (unsigned long long)(arena >> 20), (unsigned long long)(di->maxAlloc >> 20)); return 2; }
        gScratch = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, arena, NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer scratch");
        gScratchWarps = warps;
        printf("variant 5: %u persistent warps (%u x 32 work-items, work-group 32), scratch arena %llu MiB, %u units per dispatch, lazy tagged fill\n", warps, warps, (unsigned long long)(arena >> 20), units);
    }
    cur = (ServePair*)calloc(1, sizeof(ServePair));
    cur->kHashBound = dv->kHashBound; cur->kCacheFill = dv->kCacheFill; cur->kBuild = dv->kBuild; cur->prog = dv->prog;
    dv->kHashBound = dv->kCacheFill = dv->kBuild = NULL; dv->prog = NULL;
    memcpy(cur->sw, gPack.seedw, 32); memcpy(cur->kw, gPack.keyw, 32);
    if (!pairHotKernel(cur, o->packDir, perr, sizeof(perr))) { printf("FAIL: pack %s: %s\n", o->packDir, perr); return 1; }
    if (!pairBuffers(dv, di, dv->q, cur, words, gServeCacheWords, gServeSegments, o->packDir, perr, sizeof(perr))) { printf("FAIL: pack %s: %s\n", o->packDir, perr); return 1; }
    printf("pack %s: cache %.0f dataset %.0f hot %.0f check %.0f ms (%.0f ms in all); %s\n", o->packDir, cur->cacheMs, cur->datasetMs, cur->hotMs, cur->checkMs, wallMs() - t0, cur->check);
    printKernelInfo(di, cur->kHashBound, "igneum_hash_bound", (int)groupSize, "");
    dOut = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, (size_t)nonces * sizeof(uint64_t), NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer out");
    dInit = clCreateBuffer(dv->ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 32, cur->sw, &err); CL_CHECK_ERR(err, "clCreateBuffer init words");
    hOut = (uint64_t*)malloc((size_t)nonces * sizeof(uint64_t));
    strncpy(devName, di->name, 255); devName[255] = 0;
    for (k = 0; devName[k]; ++k) if (devName[k] == ' ') devName[k] = '_';
    g = gPack.persistent ? (size_t)warps * 32u : (size_t)nonces;
    for (b = -1; b < o->batches; ++b) {
        cl_uint base = (cl_uint)((uint32_t)(b + 1) * nonces);
        cl_event ev = NULL;
        double ms;
        CL_CHECK(clSetKernelArg(cur->kHashBound, 0, sizeof(cl_mem), &cur->ds));
        CL_CHECK(clSetKernelArg(cur->kHashBound, 1, sizeof(cl_mem), &dOut));
        CL_CHECK(clSetKernelArg(cur->kHashBound, 2, sizeof(cl_uint), &base));
        CL_CHECK(clSetKernelArg(cur->kHashBound, 3, sizeof(cl_uint), &mask));
        CL_CHECK(clSetKernelArg(cur->kHashBound, 4, sizeof(cl_mem), &dInit));
        if (cur->hot) CL_CHECK(clSetKernelArg(cur->kHashBound, 5, sizeof(cl_mem), &cur->hot));
        if (gPack.persistent) CL_CHECK(setScratchArgs(cur->kHashBound, cur->hot ? 6 : 5, units));
        {
            double w0 = wallMs();
            CL_CHECK(clEnqueueNDRangeKernel(dv->q, cur->kHashBound, 1, NULL, &g, &groupSize, 0, NULL, &ev));
            CL_CHECK(clWaitForEvents(1, &ev));
            ms = o->timeWall ? wallMs() - w0 : eventMs(ev);
            if (ms < 0) ms = wallMs() - w0;
            clReleaseEvent(ev);
        }
        if (b < 0) {
            warmMs = ms;
            CL_CHECK(clEnqueueReadBuffer(dv->q, dOut, CL_TRUE, 0, (size_t)nonces * sizeof(uint64_t), hOut, 0, NULL, NULL));
            fp = pf_fnv1a64((const uint32_t*)hOut, (size_t)nonces * 8u);
        } else sum += ms;
    }
    printf("warm-up dispatch (base 0): %.2f ms; %d timed dispatches of 2^%d nonces: mean %.2f ms\n", warmMs, o->batches, o->batchLog2, sum / o->batches);
    printf("RESULT pack=%s class=%s device=%s platform=%s group=%d warps=%u 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=%s\n",
           o->packDir, gPack.loadClass, devName, strcmp(di->platformName, "Apple") == 0 ? "Apple" : "other", (int)groupSize, gPack.persistent ? warps : 0u,
           gPack.persistent ? (unsigned long long)(((size_t)warps * 32u * gPack.scratchWordsPerLane * 4u) >> 20) : 0ull, cur->hotMb, cur->hotSlots, cur->hotMs, nonces, o->batches,
           cur->checked ? "PASS" : "skipped", (unsigned long long)fp, (double)nonces * (double)o->batches / (sum / 1000.0) / 1e6,
           gPack.loadsPerHash, gPack.bytesPerHash, gPack.scratchOps * 8u, o->timeWall ? "wall" : "event");
    free(hOut);
    clReleaseMemObject(dOut); clReleaseMemObject(dInit);
    if (gScratch) clReleaseMemObject(gScratch);
    releasePair(cur);
    return 0;
#endif
}

static int runServe(Device* dv, const DeviceInfo* di, const Options* o) {
#if IGNEUM_DATASET_MODE != 1
    (void)dv; (void)di; (void)o;
    printf("error 0 --serve needs a memory-hard pack (IGNEUM_DATASET_MODE 1)\n"); fflush(stdout);
    return 2;
#else
    static const uint32_t KEYW[8] = IGNEUM_KEY_INIT;
    const uint32_t words = gServeWords;
    const uint32_t mask = words - 1u;
    const uint32_t batch = 1u << (o->batchLog2 == 24 ? 22 : o->batchLog2);   /* 2^22 nonces per dispatch by default */
    size_t groupSize = 32 * (size_t)o->groupWarps;
    cl_int err = 0;
    cl_mem dOut, dInit;
    cl_uint nItems = words / 16u;
    uint64_t* hOut;
    char devName[256];
    char line[2048];
    size_t k;
    ServePair* cur;             /* the pair jobs run on */
    ServePair* prepared = NULL; /* the pair the last prepare built, until a job switches to it */
    ServePair* old = NULL;      /* the previous pair, released after the first job on the new one */
    PrepareTask* task = NULL;   /* the prepare in flight */
    if (!dv->kHashBound) { printf("error 0 the kernel source has no igneum_hash_bound (build from the pack's kernel_bound.cl, or pass --kernel)\n"); fflush(stdout); return 2; }
    cur = (ServePair*)calloc(1, sizeof(ServePair));
    cur->kHashBound = dv->kHashBound; cur->kCacheFill = dv->kCacheFill; cur->kBuild = dv->kBuild; cur->prog = dv->prog;
    dv->kHashBound = dv->kCacheFill = dv->kBuild = NULL; dv->prog = NULL;   /* owned by the pair now */
    if (gGeneric) {
        /* --pack: the first pair is the directory's pack, built and self-tested exactly like a prepared one */
        char perr[512];
        double t0 = wallMs();
        memcpy(cur->sw, gPack.seedw, 32); memcpy(cur->kw, gPack.keyw, 32);
        strncpy(cur->epochHex, gPack.epochHex, 64); cur->epochHex[64] = 0;
        strncpy(cur->dayHex, gPack.dayHex, sizeof(cur->dayHex) - 1);
        strncpy(cur->programClass, gPack.programClass, sizeof(cur->programClass) - 1); strncpy(cur->eraHex, gPack.eraHex, sizeof(cur->eraHex) - 1);
        if (!pairHotKernel(cur, o->packDir, perr, sizeof(perr))) { printf("error 0 pack %s: %s\n", o->packDir, perr); fflush(stdout); return 1; }
        if (!pairBuffers(dv, di, dv->q, cur, words, gServeCacheWords, gServeSegments, o->packDir, perr, sizeof(perr))) { printf("error 0 pack %s: %s\n", o->packDir, perr); fflush(stdout); return 1; }
        printf("info first pack %s: cache %.0f dataset %.0f check %.0f ms (%.0f ms in all); %s\n", o->packDir, cur->cacheMs, cur->datasetMs, cur->checkMs, wallMs() - t0, cur->check);
    } else {
        if (!setupCache(dv, di)) { printf("error 0 cache check failed (device cache differs from the host cache or the pack's FNV)\n"); fflush(stdout); return 1; }
        memcpy(cur->sw, SEEDW, 32); memcpy(cur->kw, KEYW, 32);
        cur->cache = gCache; gCache = NULL;
        cur->ds = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, (size_t)words * 4u, NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer dataset");
        CL_CHECK(clSetKernelArg(cur->kBuild, 0, sizeof(cl_mem), &cur->ds));
        CL_CHECK(clSetKernelArg(cur->kBuild, 1, sizeof(cl_mem), &cur->cache));
        CL_CHECK(clSetKernelArg(cur->kBuild, 2, sizeof(cl_uint), &nItems));
        countRelease(launch1D(dv, cur->kBuild, nItems, kernelMaxLocal(dv, cur->kBuild, di, 256)));
        CL_CHECK(clFinish(dv->q));
    }
    dOut = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, (size_t)batch * sizeof(uint64_t), NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer out");
    dInit = clCreateBuffer(dv->ctx, CL_MEM_READ_ONLY, 32, NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer init words");
    gMemCreated += 2;
    hOut = (uint64_t*)malloc((size_t)batch * sizeof(uint64_t));
    strncpy(devName, di->name, 255); devName[255] = 0;
    for (k = 0; devName[k]; ++k) if (devName[k] == ' ') devName[k] = '_';
    printKernelInfo(di, cur->kHashBound, "igneum_hash_bound", (int)groupSize, "info ");
    /* The select pass (5 October 2026). Before it every dispatch read back 8 bytes per nonce (16 MiB for a 2^21-nonce
     * job) over the bus and scanned them on the host; on an eGPU over USB4 that is a measurable part of every job.
     * Now a tiny kernel built here (no pack involved) writes the hits (index, hash) behind an atomic counter plus 34
     * sentinel words (the first 32 outputs, the middle and the last), and the host reads back a few hundred bytes.
     * The fault detectors read the sentinels; the found lines are printed in nonce order from the sorted hits. If a
     * chunk has more hits than the table holds (a target that loose is a test, not a block), the chunk falls back to the
     * full read. --readback full or IGNEUM_READBACK=full keeps the old path for a comparison. */
#define IG_MAX_HITS 256u
    cl_program selProg = NULL;
    cl_kernel kSelect = NULL;
    cl_mem dCount = NULL, dHits = NULL, dSentinel = NULL;
    size_t selLocal = di->maxWorkGroup < 256 ? di->maxWorkGroup : 256;
    uint32_t hCount = 0;
    uint64_t hHits[IG_MAX_HITS * 2];
    uint64_t hSentinel[34];
    unsigned long long bytesUp = 0, bytesDown = 0;
    double kernelMsSum = 0, selectMsSum = 0, readMsSum = 0, scanMsSum = 0;
    unsigned long fullFallbacks = 0;
    if (!o->readback) {
        static const char* SELECT_SRC =
            "__kernel void igneum_select(__global const ulong* out, uint n, ulong target, volatile __global uint* count,\n"
            "                            __global ulong* hits, uint maxHits, __global ulong* sentinel) {\n"
            "    uint i = (uint)get_global_id(0);\n"
            "    if (i < n) {\n"
            "        ulong h = out[i];\n"
            "        if (h <= target) { uint k = atomic_inc(count); if (k < maxHits) { hits[2u * k] = (ulong)i; hits[2u * k + 1u] = h; } }\n"
            "        if (i < 32u) sentinel[i] = h;\n"
            "        if (i == 0u) { sentinel[32] = out[n / 2u]; sentinel[33] = out[n - 1u]; }\n"
            "    }\n"
            "}\n";
        size_t selLen = strlen(SELECT_SRC);
        selProg = clCreateProgramWithSource(dv->ctx, 1, &SELECT_SRC, &selLen, &err); CL_CHECK_ERR(err, "clCreateProgramWithSource select");
        err = clBuildProgram(selProg, 1, &di->device, "-cl-std=CL1.2", NULL, NULL);
        if (err != CL_SUCCESS) {
            size_t logLen = 0; char* log;
            clGetProgramBuildInfo(selProg, di->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &logLen);
            log = (char*)calloc(logLen + 1, 1);
            if (logLen) clGetProgramBuildInfo(selProg, di->device, CL_PROGRAM_BUILD_LOG, logLen, log, NULL);
            printf("info the select pass did not build (%s): %.300s; using the full read-back\n", clErrName(err), log);
            free(log); clReleaseProgram(selProg); selProg = NULL; ((Options*)o)->readback = 1;
        } else {
            kSelect = clCreateKernel(selProg, "igneum_select", &err); CL_CHECK_ERR(err, "clCreateKernel igneum_select");
            dCount = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, 4, NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer count");
            dHits = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, IG_MAX_HITS * 2 * sizeof(uint64_t), NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer hits");
            dSentinel = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, 34 * sizeof(uint64_t), NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer sentinel");
            gMemCreated += 3;
            {
                size_t wg = 0;
                if (clGetKernelWorkGroupInfo(kSelect, di->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(wg), &wg, NULL) == CL_SUCCESS && wg && wg < selLocal) selLocal = wg;
            }
        }
    }
    printf("info readback %s (per dispatch of %u nonces: %s)\n", o->readback ? "full" : "select",
           batch, o->readback ? "8 bytes per nonce come back and the host scans them" : "the hits and 34 sentinel words come back; the GPU scans");
    printf("ready opencl %s platform %s pack %s dataset-log2 %d batch %u exchange %d prepare %d path %s\n", devName, di->platformName,
           gGeneric ? gPack.seedString : IGNEUM_SEED_STRING, log2u32(words), batch, dv->exchange, o->noPrepare ? 0 : 1, gGeneric ? "prebuilt-generic" : "compiled-in");
    fflush(stdout);

    /* Fault detection (4 October 2026, after the gfx1036 run that completed 2,000 jobs a second with no hash after
     * 600 s): any OpenCL error in the job path is fatal (exit 3, the miner restarts the worker), the dispatch event
     * must report CL_COMPLETE, a chunk that runs 20x faster per nonce than the running mean is a fault, and two
     * consecutive chunks with the same output signature (different nonces or init words give different outputs) mean
     * the kernel did not run. A stats line every 200 jobs shows the live event and buffer counts. */
    long gFaultTestChunk = getenv("IGNEUM_FAULT_TEST") ? atol(getenv("IGNEUM_FAULT_TEST")) : -1;
    double meanNsPerNonce = 0;      /* running mean of wall ns per nonce over the chunks so far */
    unsigned long chunksSeen = 0, jobsSeen = 0;
    unsigned long long noncesSeen = 0;      /* for the mean chunk in the stats line */
    uint64_t prevSig = 0; int havePrevSig = 0;
    double jobMsSum = 0;
#define SERVE_FATAL(jobIdStr, fmt, ...) do { printf("error %s worker fault: " fmt "; exiting 3 so the miner restarts the worker\n", jobIdStr, __VA_ARGS__); fflush(stdout); exit(3); } while (0)
#define SERVE_CHECK(jobIdStr, call) do { cl_int e_ = (call); if (e_ != CL_SUCCESS) SERVE_FATAL(jobIdStr, "%s failed with %s (%d)", #call, clErrName(e_), (int)e_); } while (0)
    while (fgets(line, sizeof(line), stdin)) {
        char* f[12];                /* job + 7 fields, plus the class= and era= tokens of a class v3 line */
        int nf = 0;
        char* tok;
        char* save = NULL;
        char jobId[64];
        char wantClass[8], wantEra[65];
        uint8_t prehash[32], epochSeed[32], daySeed[256];
        size_t prehashLen = 0, epochLen = 0, dayLen = 0;
        unsigned long long target = 0, nonceStart = 0, nonceCount = 0;
        uint32_t sw[8], kw[8];
        uint64_t remaining, hashes = 0;
        uint32_t hi, lo;
        double t0;
        int failed = 0, switched = 0;
        line[strcspn(line, "\r\n")] = 0;
        /* A finished prepare is reported here, between lines (the thread never prints) */
        if (task && task->done) {
            if (task->result) {
                ServePair* p = task->result;
                {
                    uint8_t eb[32], db[256]; size_t el = 0, dl = 0;
                    if (unhexBuf(p->epochHex, eb, 32, &el) && el == 32) seedWordsFromBytes(eb, 32, p->sw);
                    if (unhexBuf(p->dayHex, db, sizeof(db), &dl)) seedWordsFromBytes(db, dl, p->kw);
                }
                if (prepared) releasePair(prepared);
                prepared = p;
                printf("prepared %s %s %.1f build %.1f cache %.1f dataset %.1f check %.1f resident 2 programs 2 datasets; %s\n", p->epochHex, p->dayHex, task->doneAt - task->t0, p->buildMs, p->cacheMs, p->datasetMs, p->checkMs, p->check);
            } else {
                printf("prepare-failed %s %s %s\n", task->epochHex, task->dayHex, task->error);
            }
            fflush(stdout);
            free(task); task = NULL;
        }
        for (tok = strtok_r(line, " ", &save); tok && nf < 12; tok = strtok_r(NULL, " ", &save)) f[nf++] = tok;
        if (nf == 0) continue;
        if (strcmp(f[0], "quit") == 0) break;
        takeClassTokens(f, &nf, wantClass, sizeof(wantClass), wantEra, sizeof(wantEra));
        if (strcmp(f[0], "prepare") == 0) {
            if (o->noPrepare) { printf("info ignored (started with --no-prepare): prepare\n"); fflush(stdout); continue; }
            if (nf < 4) { printf("prepare-failed %s %s this ahead-of-time worker needs a pack directory as the third field (igneum-miner --prepare-packs <dir>)\n", nf > 1 ? f[1] : "0", nf > 2 ? f[2] : "0"); fflush(stdout); continue; }
            if (strlen(f[1]) != 64 || strlen(f[2]) >= 500) { printf("prepare-failed %s %s bad field (epoch_seed 64 hex, day_seed hex)\n", f[1], f[2]); fflush(stdout); continue; }
            if (task) { printf("prepare-failed %s %s a prepare is still running\n", f[1], f[2]); fflush(stdout); continue; }
            if (prepared && strcmp(prepared->epochHex, f[1]) == 0 && strcmp(prepared->dayHex, f[2]) == 0) { printf("prepared %s %s 0 (already resident)\n", f[1], f[2]); fflush(stdout); continue; }
            task = (PrepareTask*)calloc(1, sizeof(PrepareTask));
            task->dv = dv; task->di = di; task->words = words; task->cacheWords = gServeCacheWords; task->segments = gServeSegments; task->t0 = wallMs();
            strncpy(task->epochHex, f[1], 64); strncpy(task->dayHex, f[2], sizeof(task->dayHex) - 1); strncpy(task->packDir, f[3], sizeof(task->packDir) - 1);
            strncpy(task->wantClass, wantClass, sizeof(task->wantClass) - 1); strncpy(task->wantEra, wantEra, sizeof(task->wantEra) - 1);
            if (!startPrepareThread(task)) { printf("prepare-failed %s %s cannot start the prepare thread\n", f[1], f[2]); fflush(stdout); free(task); task = NULL; continue; }
            printf("info prepare started for epoch %.16s day %s from %s (builds in the background)\n", f[1], f[2], f[3]); fflush(stdout);
            continue;
        }
        if (strcmp(f[0], "job") != 0) { printf("info ignored line\n"); fflush(stdout); continue; }
        strncpy(jobId, nf > 1 ? f[1] : "0", 63); jobId[63] = 0;
        if (nf < 8) { printf("error %s malformed job line (need 7 fields after job)\n", jobId); fflush(stdout); continue; }
        if (!unhexBuf(f[2], prehash, 32, &prehashLen) || prehashLen != 32 || sscanf(f[3], "%llx", &target) != 1 ||
            sscanf(f[4], "%llu", &nonceStart) != 1 || sscanf(f[5], "%llu", &nonceCount) != 1 ||
            !unhexBuf(f[6], epochSeed, 32, &epochLen) || epochLen != 32 || !unhexBuf(f[7], daySeed, sizeof(daySeed), &dayLen)) {
            printf("error %s bad field (prehash 64 hex, target 16 hex, nonce_start u64, nonce_count u64, epoch_seed 64 hex, day_seed hex)\n", jobId); fflush(stdout); continue;
        }
        if (nonceCount == 0 || nonceCount % 32 != 0 || (nonceStart & 31) != 0) { printf("error %s nonce_start must be 32-aligned and nonce_count a non-zero multiple of 32\n", jobId); fflush(stdout); continue; }
        seedWordsFromBytes(epochSeed, 32, sw);
        seedWordsFromBytes(daySeed, dayLen, kw);
        t0 = wallMs();
        if (!pairIsClass(cur, f[6], f[7], sw, kw, wantClass, wantEra)) {
            char why[256];
            if (pairIsClass(prepared, f[6], f[7], sw, kw, wantClass, wantEra)) {
                /* The prepared pair: switch now, release the old one after this job */
                if (old) releasePair(old);
                old = cur; cur = prepared; prepared = NULL; switched = 1;
                printf("info switched to the prepared pair epoch %.16s day %s (class %s) in %.2f ms\n", cur->epochHex, cur->dayHex, cur->programClass[0] ? cur->programClass : "v2", wallMs() - t0); fflush(stdout);
            } else if (pairIs(cur, f[6], f[7], sw, kw) && !pf_pack_class_ok(cur->programClass[0] ? cur->programClass : "v2", cur->eraHex, wantClass, wantEra, why, sizeof(why))) {
                /* Counter ASIC 2.0: right seeds, wrong class or era; the miner prepares the pair from a pack of the
                 * class the chain is on (the need line), and this pack is never mined */
                printf("need %s %s\n", f[6], f[7]);
                printf("error %s pack %s: %s\n", jobId, o->packDir ? o->packDir : "(compiled-in)", why);
                fflush(stdout); continue;
            } else if (cur->epochHex[0] ? !hexEq(cur->epochHex, f[6]) : memcmp(sw, cur->sw, 32) != 0) {
                printf("need %s %s\n", f[6], f[7]);   /* the miner prepares this pair (4 October 2026) */
                printf("error %s epoch seed mismatch: this worker holds %s%s (program words %08x %08x ...)%s, the job is for epoch %.16s (bare seed words %08x %08x ...); send prepare with a pack directory, or run igneum-miner export-pack and rebuild\n",
                       jobId, cur->epochHex[0] ? "prepared epoch " : "pack \"" IGNEUM_SEED_STRING "\"", cur->epochHex[0] ? cur->epochHex : "", cur->sw[0], cur->sw[1], prepared ? " plus one prepared pair" : "", f[6], sw[0], sw[1]);
                fflush(stdout); continue;
            } else {
                printf("need %s %s\n", f[6], f[7]);
                printf("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, or run igneum-miner export-pack and rebuild\n",
                       jobId, cur->kw[0], cur->kw[1], f[7], kw[0], kw[1]);
                fflush(stdout); continue;
            }
        }
        remaining = nonceCount;
        hi = (uint32_t)(nonceStart >> 32); lo = (uint32_t)nonceStart;
        while (remaining > 0) {
            uint64_t room = (uint64_t)(0xffffffffu - lo) + 1ull;
            uint64_t chunk64 = remaining < batch ? remaining : batch;
            uint32_t chunk, i;
            uint32_t iw[8];
            uint8_t b[49];
            cl_uint baseNonce, maskArg = mask;
            cl_event ev;
            if (chunk64 > room) chunk64 = room;
            chunk = (uint32_t)chunk64;
            memcpy(b, "igneum-block/", 13);
            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);
            seedWordsFromBytes(b, 49, iw);
            {
            double c0 = wallMs(), cms, kernelMs = -1.0, selectMs = 0.0, readMs = 0.0, scanMs = 0.0, r0;
            cl_int status = 0;
            size_t g = ((chunk + groupSize - 1) / groupSize) * groupSize;
            uint64_t sig;
            int useSelect = (kSelect != NULL);
            uint32_t nHits = 0;
            SERVE_CHECK(jobId, clEnqueueWriteBuffer(dv->q, dInit, CL_TRUE, 0, 32, iw, 0, NULL, NULL));
            bytesUp += 32;
            if (useSelect) { hCount = 0; SERVE_CHECK(jobId, clEnqueueWriteBuffer(dv->q, dCount, CL_TRUE, 0, 4, &hCount, 0, NULL, NULL)); bytesUp += 4; }
            baseNonce = lo;
            SERVE_CHECK(jobId, clSetKernelArg(cur->kHashBound, 0, sizeof(cl_mem), &cur->ds));
            SERVE_CHECK(jobId, clSetKernelArg(cur->kHashBound, 1, sizeof(cl_mem), &dOut));
            SERVE_CHECK(jobId, clSetKernelArg(cur->kHashBound, 2, sizeof(cl_uint), &baseNonce));
            SERVE_CHECK(jobId, clSetKernelArg(cur->kHashBound, 3, sizeof(cl_uint), &maskArg));
            SERVE_CHECK(jobId, clSetKernelArg(cur->kHashBound, 4, sizeof(cl_mem), &dInit));
            if (cur->hot) SERVE_CHECK(jobId, clSetKernelArg(cur->kHashBound, 5, sizeof(cl_mem), &cur->hot));
            ev = NULL;
            if (gFaultTestChunk >= 0 && (long)chunksSeen >= gFaultTestChunk) {
                /* IGNEUM_FAULT_TEST=N (test only): from chunk N on, behave like a runtime that answers every call with
                 * success and runs nothing, so the detectors below are exercised on a healthy device. */
            } else {
            SERVE_CHECK(jobId, clEnqueueNDRangeKernel(dv->q, cur->kHashBound, 1, NULL, &g, &groupSize, 0, NULL, &ev));
            ++gEvCreated;
            /* Wait on the dispatch event, not clFinish: the runtime can sleep the thread on an event, where clFinish
             * on some drivers spins one core for the whole dispatch. The event must then say CL_COMPLETE: a runtime
             * that has dropped its device answers the wait at once with a negative status. */
            err = clWaitForEvents(1, &ev);
            if (err != CL_SUCCESS) { countRelease(ev); SERVE_FATAL(jobId, "clWaitForEvents on the dispatch returned %s (%d)", clErrName(err), (int)err); }
            err = clGetEventInfo(ev, CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(status), &status, NULL);
            kernelMs = eventMs(ev);
            countRelease(ev);
            if (err != CL_SUCCESS) SERVE_FATAL(jobId, "clGetEventInfo on the dispatch returned %s (%d)", clErrName(err), (int)err);
            if (status != CL_COMPLETE) SERVE_FATAL(jobId, "the dispatch event ended with status %d, not CL_COMPLETE (a device reset or a lost context)", (int)status);
            if (useSelect) {
                cl_event evs = NULL;
                cl_uint nArg = chunk, maxArg = IG_MAX_HITS;
                cl_ulong tArg = (cl_ulong)target;
                size_t gs = ((chunk + selLocal - 1) / selLocal) * selLocal;
                SERVE_CHECK(jobId, clSetKernelArg(kSelect, 0, sizeof(cl_mem), &dOut));
                SERVE_CHECK(jobId, clSetKernelArg(kSelect, 1, sizeof(cl_uint), &nArg));
                SERVE_CHECK(jobId, clSetKernelArg(kSelect, 2, sizeof(cl_ulong), &tArg));
                SERVE_CHECK(jobId, clSetKernelArg(kSelect, 3, sizeof(cl_mem), &dCount));
                SERVE_CHECK(jobId, clSetKernelArg(kSelect, 4, sizeof(cl_mem), &dHits));
                SERVE_CHECK(jobId, clSetKernelArg(kSelect, 5, sizeof(cl_uint), &maxArg));
                SERVE_CHECK(jobId, clSetKernelArg(kSelect, 6, sizeof(cl_mem), &dSentinel));
                SERVE_CHECK(jobId, clEnqueueNDRangeKernel(dv->q, kSelect, 1, NULL, &gs, &selLocal, 0, NULL, &evs));
                ++gEvCreated;
                err = clWaitForEvents(1, &evs);
                if (err != CL_SUCCESS) { countRelease(evs); SERVE_FATAL(jobId, "clWaitForEvents on the select pass returned %s (%d)", clErrName(err), (int)err); }
                selectMs = eventMs(evs);
                countRelease(evs);
            }
            }
            r0 = wallMs();
            if (useSelect) {
                SERVE_CHECK(jobId, clEnqueueReadBuffer(dv->q, dCount, CL_TRUE, 0, 4, &hCount, 0, NULL, NULL));
                bytesDown += 4;
                nHits = hCount;
                if (nHits > IG_MAX_HITS) {
                    /* more hits than the table holds: this chunk takes the full path (the test target case) */
                    ++fullFallbacks;
                    useSelect = 0;
                } else {
                    if (nHits) { SERVE_CHECK(jobId, clEnqueueReadBuffer(dv->q, dHits, CL_TRUE, 0, (size_t)nHits * 2 * sizeof(uint64_t), hHits, 0, NULL, NULL)); bytesDown += (unsigned long long)nHits * 16; }
                    SERVE_CHECK(jobId, clEnqueueReadBuffer(dv->q, dSentinel, CL_TRUE, 0, 34 * sizeof(uint64_t), hSentinel, 0, NULL, NULL));
                    bytesDown += 34 * 8;
                }
            }
            if (!useSelect) {
                SERVE_CHECK(jobId, clEnqueueReadBuffer(dv->q, dOut, CL_TRUE, 0, (size_t)chunk * sizeof(uint64_t), hOut, 0, NULL, NULL));
                bytesDown += (unsigned long long)chunk * 8;
            }
            readMs = wallMs() - r0;
            cms = wallMs() - c0;
            /* Plausibility: wall time per nonce against the running mean (the first chunk sets it; a chunk is a full
             * batch except at the 32-bit boundary, so per nonce is the comparable unit). 20x faster = the kernel did not run. */
            {
                double ns = cms * 1e6 / (double)chunk;
                if (chunksSeen >= 4 && meanNsPerNonce > 0 && ns * 20.0 < meanNsPerNonce)
                    SERVE_FATAL(jobId, "%u nonces reported complete in %.3f ms, %.0fx faster than the running mean of %.2f ms per million (the runtime is not running the kernel)", chunk, cms, meanNsPerNonce / ns, meanNsPerNonce / 1e3);
                meanNsPerNonce = chunksSeen == 0 ? ns : meanNsPerNonce + (ns - meanNsPerNonce) / (double)(chunksSeen + 1 < 64 ? chunksSeen + 1 : 64);
                ++chunksSeen; noncesSeen += chunk;
            }
            /* Stale output: the first, middle and last words plus an FNV of the first 32. Two chunks never share it.
             * On the select path the same 34 words come from the sentinel buffer the select pass wrote. */
            if (useSelect) sig = fnv1a64(hSentinel, 32 * sizeof(uint64_t)) ^ hSentinel[32] ^ hSentinel[33];
            else sig = fnv1a64(hOut, 32 * sizeof(uint64_t)) ^ hOut[chunk / 2] ^ hOut[chunk - 1];
            if (chunk >= 64) {
                if (havePrevSig && sig == prevSig) SERVE_FATAL(jobId, "the output buffer is unchanged since the previous dispatch (signature %016llx): the kernel did not run", (unsigned long long)sig);
                prevSig = sig; havePrevSig = 1;
            }
            r0 = wallMs();
            if (useSelect) {
                /* nonce order, as the full scan printed them (the miner submits the first found of a job) */
                uint32_t a, b;
                for (a = 1; a < nHits; ++a) {
                    uint64_t ki = hHits[2 * a], kh = hHits[2 * a + 1];
                    for (b = a; b > 0 && hHits[2 * (b - 1)] > ki; --b) { hHits[2 * b] = hHits[2 * (b - 1)]; hHits[2 * b + 1] = hHits[2 * (b - 1) + 1]; }
                    hHits[2 * b] = ki; hHits[2 * b + 1] = kh;
                }
                for (a = 0; a < nHits; ++a) {
                    uint32_t idx = (uint32_t)hHits[2 * a];
                    unsigned long long nonce = ((unsigned long long)hi << 32) | (unsigned long long)(uint32_t)(lo + idx);
                    printf("found %s %llu %016llx\n", jobId, nonce, (unsigned long long)hHits[2 * a + 1]);
                }
            } else {
                for (i = 0; i < chunk; ++i) if (hOut[i] <= target) {
                    unsigned long long nonce = ((unsigned long long)hi << 32) | (unsigned long long)(uint32_t)(lo + i);
                    printf("found %s %llu %016llx\n", jobId, nonce, (unsigned long long)hOut[i]);
                }
            }
            scanMs = wallMs() - r0;
            if (kernelMs > 0) kernelMsSum += kernelMs;
            selectMsSum += selectMs; readMsSum += readMs; scanMsSum += scanMs;
            }
            fflush(stdout);
            hashes += chunk;
            remaining -= chunk;
            if (chunk64 == room) { hi += 1u; lo = 0u; } else lo += chunk;
        }
        if (failed) continue;
        printf("done %s %llu %.2f\n", jobId, (unsigned long long)hashes, wallMs() - t0);
        jobMsSum += wallMs() - t0;
        ++jobsSeen;
        if (jobsSeen % 200 == 0) {
            printf("info stats jobs %lu chunks %lu mean job %.1f ms mean %.2f ms per million nonces; events created %lu released %lu live %lu; buffers created %lu released %lu live %lu\n",
                   jobsSeen, chunksSeen, jobMsSum / (double)jobsSeen, meanNsPerNonce / 1e3, gEvCreated, gEvReleased, gEvCreated - gEvReleased, gMemCreated, gMemReleased, gMemCreated - gMemReleased);
            /* The transfer and time budget per chunk (5 October 2026): bytes up (init words, the count reset), bytes
             * down (hits and sentinels, or 8 bytes per nonce on the full path), and the mean device time of the hash
             * kernel (event profiling), the select pass, the blocking read-back and the host scan. */
            printf("info transfers per chunk: up %.0f B, down %.0f B (%s, %lu full fallbacks); mean per chunk: kernel %.2f ms (device), select %.3f ms, read-back %.2f ms, scan %.2f ms, chunk wall %.2f ms\n",
                   (double)bytesUp / (double)chunksSeen, (double)bytesDown / (double)chunksSeen, kSelect ? "select" : "full", fullFallbacks,
                   kernelMsSum / (double)chunksSeen, selectMsSum / (double)chunksSeen, readMsSum / (double)chunksSeen, scanMsSum / (double)chunksSeen, meanNsPerNonce * ((double)noncesSeen / (double)chunksSeen) / 1e6);
            if (gEvCreated - gEvReleased > 16 || gMemCreated - gMemReleased > 12) SERVE_FATAL(jobId, "object leak: %lu events and %lu buffers live after %lu jobs", gEvCreated - gEvReleased, gMemCreated - gMemReleased, jobsSeen);
        }
        if (switched && old) { releasePair(old); old = NULL; printf("info dropped the previous pair (its program, cache and dataset)\n"); }
        fflush(stdout);
    }
    free(hOut);
    clReleaseMemObject(dInit);
    clReleaseMemObject(dOut);
    gMemReleased += 2;
    if (kSelect) { clReleaseKernel(kSelect); clReleaseMemObject(dCount); clReleaseMemObject(dHits); clReleaseMemObject(dSentinel); gMemReleased += 3; }
    if (selProg) clReleaseProgram(selProg);
    if (chunksSeen) printf("info transfers total: %lu chunks, up %llu B, down %llu B (%s, %lu full fallbacks)\n", chunksSeen, bytesUp, bytesDown, kSelect ? "select" : "full", fullFallbacks);
    if (old) releasePair(old);
    if (prepared) releasePair(prepared);
    releasePair(cur);
    return 0;
#endif
}

/* ---------------------------------------------------------------------------------------------
 * --memprobe (5 October 2026): what the device itself can do with the access pattern of the hash, with no pack and no
 * program. Three kernels built from the text below:
 *   chase  one dependent random 4-byte load per step per lane (the next address comes from the loaded word), so the
 *          time per step at a small lane count is the loaded-latency of one random read, and the loads per second at a
 *          large lane count is the device's random-read throughput for a dependent chain (the hash is 128 of these)
 *   indep  eight independent chains per lane: the throughput when latency is hidden inside one lane
 *   alu    an integer multiply-add-rotate chain with no memory: the achieved integer rate, which moves with the clock
 * Sizes 4 MiB (cache-resident), 64 MiB (the last-level cache on RDNA 3 and 4 is 64 MiB or more, approximate) and
 * 1024 MiB (the dataset size: DRAM). Times from event profiling (wall on Apple). Every number is printed with the
 * configuration that produced it. Rates are G loads/s = 1e9 loads per second; ns per load = time / steps.
 */
static const char* PROBE_SRC =
    "static inline uint pm_mix(uint x) { x ^= x >> 16; x *= 0x7feb352du; x ^= x >> 15; x *= 0x846ca68bu; x ^= x >> 16; return x; }\n"
    "__kernel void probe_fill(__global uint* ds, uint n) { uint i = (uint)get_global_id(0); if (i < n) ds[i] = pm_mix(i ^ 0x9E3779B9u); }\n"
    "__kernel void probe_chase(__global const uint* ds, uint mask, uint steps, uint seed, __global uint* out) {\n"
    "    uint x = pm_mix((uint)get_global_id(0) ^ seed);\n"
    "    for (uint s = 0u; s < steps; ++s) x = ds[x & mask] ^ (x * 0x9E3779B1u + s);\n"
    "    out[get_global_id(0)] = x;\n"
    "}\n"
    "__kernel void probe_indep(__global const uint* ds, uint mask, uint steps, uint seed, __global uint* out) {\n"
    "    uint g = (uint)get_global_id(0);\n"
    "    uint 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"
    "    uint 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 (uint 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"
    "__kernel void probe_line16(__global const uint4* ds, uint vecMask, uint steps, uint seed, __global uint* out) {\n"
    "    uint x = pm_mix((uint)get_global_id(0) ^ seed);\n"
    "    for (uint s = 0u; s < steps; ++s) { uint4 a = ds[x & vecMask]; x = (a.x ^ a.y ^ a.z ^ a.w) ^ (x * 0x9E3779B1u + s); }\n"
    "    out[get_global_id(0)] = x;\n"
    "}\n"
    "__kernel void probe_line(__global const uint4* ds, uint lineMask, uint steps, uint seed, __global uint* out) {\n"
    "    uint x = pm_mix((uint)get_global_id(0) ^ seed);\n"
    "    for (uint s = 0u; s < steps; ++s) {\n"
    "        uint l = (x & lineMask) * 4u;\n"
    "        uint4 a = ds[l], b = ds[l + 1u], c = ds[l + 2u], d = ds[l + 3u];\n"
    "        x = (a.x ^ b.y ^ c.z ^ d.w) ^ (x * 0x9E3779B1u + s);\n"
    "    }\n"
    "    out[get_global_id(0)] = x;\n"
    "}\n"
    "__kernel void probe_stream(__global const uint4* ds, uint perLane, __global uint* out) {\n"
    "    uint g = (uint)get_global_id(0), n = (uint)get_global_size(0);\n"
    "    uint4 acc = (uint4)(0u, 0u, 0u, 0u);\n"
    "    for (uint s = 0u; s < perLane; ++s) acc ^= ds[s * n + g];\n"
    "    out[g] = acc.x ^ acc.y ^ acc.z ^ acc.w;\n"
    "}\n"
    "__kernel void probe_alu(uint steps, uint seed, __global uint* out) {\n"
    "    uint g = (uint)get_global_id(0); uint x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u;\n"
    "    for (uint s = 0u; s < steps; ++s) { x = x * 0x9E3779B1u + rotate(y, 7u); y = (y ^ x) + s; }\n"
    "    out[g] = x ^ y;\n"
    "}\n";

/* One timed launch, best of `reps`, in ms (event time, or wall when the platform's events are unusable). */
static double probeLaunch(Device* dv, const Options* o, cl_kernel k, size_t global, size_t local, int reps, int seedArg, cl_uint seed) {
    double best = -1.0;
    int r;
    for (r = 0; r < reps; ++r) {
        cl_event ev = NULL;
        double w0 = wallMs(), ms;
        /* a fresh seed per repetition: a replay of the same addresses would be served from the last-level cache
         * (65,536 reads x 64 B = 4 MiB fits any of them) and read as DRAM latency (seen on the 9070 XT, 5 October) */
        if (seedArg >= 0) { cl_uint s = seed + (cl_uint)r * 0x9E3779B9u; CL_CHECK(clSetKernelArg(k, (cl_uint)seedArg, sizeof(cl_uint), &s)); }
        CL_CHECK(clEnqueueNDRangeKernel(dv->q, k, 1, NULL, &global, &local, 0, NULL, &ev));
        CL_CHECK(clWaitForEvents(1, &ev));
        ms = o->timeWall ? wallMs() - w0 : eventMs(ev);
        if (ms < 0) ms = wallMs() - w0;
        clReleaseEvent(ev);
        if (best < 0 || ms < best) best = ms;
    }
    return best;
}

static int runMemprobe(Device* dv, const DeviceInfo* di, const Options* o) {
    cl_int err = 0;
    cl_program prog;
    cl_kernel kFill, kChase, kIndep, kAlu, kLine, kStream, kLine16;
    size_t srcLen = strlen(PROBE_SRC);
    int sizes[3] = { 4, 64, 1024 }, nSizes = 3, si;
    size_t lanesList[8] = { 256, 1024, 1u << 12, 1u << 14, 1u << 16, 1u << 18, 1u << 20, 1u << 22 };
    const int nLanes = 8;
    size_t groups[2] = { 32, 256 };
    const cl_uint STEPS = 256u, ALU_STEPS = 4096u;
    const size_t maxLanes = 1u << 22;
    cl_mem dOut;
    if (o->probeMib > 0) { sizes[0] = o->probeMib; nSizes = 1; }
    prog = clCreateProgramWithSource(dv->ctx, 1, &PROBE_SRC, &srcLen, &err); CL_CHECK_ERR(err, "clCreateProgramWithSource probe");
    err = clBuildProgram(prog, 1, &di->device, "-cl-std=CL1.2", NULL, NULL);
    if (err != CL_SUCCESS) {
        size_t logLen = 0; char* log;
        clGetProgramBuildInfo(prog, di->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &logLen);
        log = (char*)calloc(logLen + 1, 1);
        if (logLen) clGetProgramBuildInfo(prog, di->device, CL_PROGRAM_BUILD_LOG, logLen, log, NULL);
        printf("memprobe: build FAILED (%s)\n--- build log ---\n%s\n--- end of build log ---\n", clErrName(err), log);
        free(log);
        return 2;
    }
    kFill = clCreateKernel(prog, "probe_fill", &err); CL_CHECK_ERR(err, "probe_fill");
    kChase = clCreateKernel(prog, "probe_chase", &err); CL_CHECK_ERR(err, "probe_chase");
    kIndep = clCreateKernel(prog, "probe_indep", &err); CL_CHECK_ERR(err, "probe_indep");
    kAlu = clCreateKernel(prog, "probe_alu", &err); CL_CHECK_ERR(err, "probe_alu");
    kLine = clCreateKernel(prog, "probe_line", &err); CL_CHECK_ERR(err, "probe_line");
    kLine16 = clCreateKernel(prog, "probe_line16", &err); CL_CHECK_ERR(err, "probe_line16");
    kStream = clCreateKernel(prog, "probe_stream", &err); CL_CHECK_ERR(err, "probe_stream");
    dOut = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, maxLanes * 4u, NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer probe out");
    printf("memprobe on [%s] %s, driver %s, %u compute units, %u MHz, %s time\n", di->platformName, di->name, di->driver, di->computeUnits, di->clockMHz, o->timeWall ? "wall" : "device event");
    printKernelInfo(di, kChase, "probe_chase", 32, "memprobe ");
    printKernelInfo(di, kChase, "probe_chase", 256, "memprobe ");
    printf("| probe | MiB | work-group | lanes in flight | steps per lane | best ms | G loads/s | ns per dependent load |\n|---|---|---|---|---|---|---|---|\n");
    for (si = 0; si < nSizes; ++si) {
        int mib = sizes[si];
        uint64_t bytes = (uint64_t)mib << 20;
        cl_uint words = (cl_uint)(bytes / 4ull), mask = words - 1u, n = words;
        cl_mem dDs;
        size_t gi, li;
        size_t fillLocal = di->maxWorkGroup < 256 ? di->maxWorkGroup : 256;
        size_t fillGlobal = ((size_t)words + fillLocal - 1) / fillLocal * fillLocal;
        if ((uint64_t)di->maxAlloc < bytes) { printf("| chase | %d | skipped: max alloc %llu MiB | | | | | |\n", mib, (unsigned long long)(di->maxAlloc >> 20)); continue; }
        dDs = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, (size_t)bytes, NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer probe dataset");
        CL_CHECK(clSetKernelArg(kFill, 0, sizeof(cl_mem), &dDs));
        CL_CHECK(clSetKernelArg(kFill, 1, sizeof(cl_uint), &n));
        probeLaunch(dv, o, kFill, fillGlobal, fillLocal, 1, -1, 0u);
        for (gi = 0; gi < 2; ++gi) {
            size_t local = groups[gi];
            if (local > di->maxWorkGroup) continue;
            for (li = 0; li < (size_t)nLanes; ++li) {
                size_t lanes = lanesList[li];
                cl_uint seed = (cl_uint)(0x1234567u + (cl_uint)li * 977u);
                double ms;
                if (lanes < local) continue;
                CL_CHECK(clSetKernelArg(kChase, 0, sizeof(cl_mem), &dDs));
                CL_CHECK(clSetKernelArg(kChase, 1, sizeof(cl_uint), &mask));
                CL_CHECK(clSetKernelArg(kChase, 2, sizeof(cl_uint), &STEPS));
                CL_CHECK(clSetKernelArg(kChase, 3, sizeof(cl_uint), &seed));
                CL_CHECK(clSetKernelArg(kChase, 4, sizeof(cl_mem), &dOut));
                ms = probeLaunch(dv, o, kChase, lanes, local, 3, 3, seed);
                printf("| chase | %d | %llu | %llu | %u | %.3f | %.3f | %.0f |\n", mib, (unsigned long long)local, (unsigned long long)lanes, STEPS, ms,
                       (double)lanes * (double)STEPS / (ms / 1000.0) / 1e9, ms * 1e6 / (double)STEPS);
                fflush(stdout);
            }
        }
        {
            size_t local = di->maxWorkGroup < 256 ? di->maxWorkGroup : 256;
            size_t lanes;
            for (lanes = 1u << 16; lanes <= maxLanes; lanes <<= 2) {
                cl_uint seed = 0x7654321u;
                double ms;
                CL_CHECK(clSetKernelArg(kIndep, 0, sizeof(cl_mem), &dDs));
                CL_CHECK(clSetKernelArg(kIndep, 1, sizeof(cl_uint), &mask));
                CL_CHECK(clSetKernelArg(kIndep, 2, sizeof(cl_uint), &STEPS));
                CL_CHECK(clSetKernelArg(kIndep, 3, sizeof(cl_uint), &seed));
                CL_CHECK(clSetKernelArg(kIndep, 4, sizeof(cl_mem), &dOut));
                ms = probeLaunch(dv, o, kIndep, lanes, local, 3, 3, seed);
                printf("| indep x8 | %d | %llu | %llu | %u | %.3f | %.3f | (8 loads in flight per lane) |\n", mib, (unsigned long long)local, (unsigned long long)lanes, STEPS, ms,
                       (double)lanes * 8.0 * (double)STEPS / (ms / 1000.0) / 1e9);
                fflush(stdout);
            }
        }
        {
            /* Random 16-byte reads (one uint4) in a dependent chain: the W = 16 width of the read-width experiment. */
            size_t local = di->maxWorkGroup < 256 ? di->maxWorkGroup : 256;
            size_t lanes;
            cl_uint vecMask = (words / 4u) - 1u;
            for (lanes = 1u << 14; lanes <= maxLanes; lanes <<= 2) {
                cl_uint seed = 0x2718281u;
                double ms;
                CL_CHECK(clSetKernelArg(kLine16, 0, sizeof(cl_mem), &dDs));
                CL_CHECK(clSetKernelArg(kLine16, 1, sizeof(cl_uint), &vecMask));
                CL_CHECK(clSetKernelArg(kLine16, 2, sizeof(cl_uint), &STEPS));
                CL_CHECK(clSetKernelArg(kLine16, 3, sizeof(cl_uint), &seed));
                CL_CHECK(clSetKernelArg(kLine16, 4, sizeof(cl_mem), &dOut));
                ms = probeLaunch(dv, o, kLine16, lanes, local, 3, 3, seed);
                printf("| line 16 B | %d | %llu | %llu | %u | %.3f | %.3f G reads/s | %.1f GB/s in 16 B reads |\n", mib, (unsigned long long)local, (unsigned long long)lanes, STEPS, ms,
                       (double)lanes * (double)STEPS / (ms / 1000.0) / 1e9, (double)lanes * (double)STEPS * 16.0 / (ms / 1000.0) / 1e9);
                fflush(stdout);
            }
        }
        {
            /* Random 64-byte lines (16 words, four uint4 loads) in a dependent chain: lines per second against the
             * 4-byte chase above says what one random 4-byte read costs the memory system. If the two rates are
             * equal, every 4-byte read fetches a whole line; if lines/s is a quarter of loads/s, reads cost a 16-byte
             * sector. */
            size_t local = di->maxWorkGroup < 256 ? di->maxWorkGroup : 256;
            size_t lanes;
            cl_uint lineMask = (words / 16u) - 1u;
            for (lanes = 1u << 14; lanes <= maxLanes; lanes <<= 2) {
                cl_uint seed = 0x3141592u;
                double ms;
                CL_CHECK(clSetKernelArg(kLine, 0, sizeof(cl_mem), &dDs));
                CL_CHECK(clSetKernelArg(kLine, 1, sizeof(cl_uint), &lineMask));
                CL_CHECK(clSetKernelArg(kLine, 2, sizeof(cl_uint), &STEPS));
                CL_CHECK(clSetKernelArg(kLine, 3, sizeof(cl_uint), &seed));
                CL_CHECK(clSetKernelArg(kLine, 4, sizeof(cl_mem), &dOut));
                ms = probeLaunch(dv, o, kLine, lanes, local, 3, 3, seed);
                printf("| line 64 B | %d | %llu | %llu | %u | %.3f | %.3f G lines/s | %.1f GB/s in lines |\n", mib, (unsigned long long)local, (unsigned long long)lanes, STEPS, ms,
                       (double)lanes * (double)STEPS / (ms / 1000.0) / 1e9, (double)lanes * (double)STEPS * 64.0 / (ms / 1000.0) / 1e9);
                fflush(stdout);
            }
        }
        {
            /* Coalesced read of the whole buffer (uint4 per lane per step, consecutive lanes consecutive addresses):
             * the sequential bandwidth. Against the card's rated figure this says whether the memory clock is in its
             * full state; a card parked in a middle memory state shows about half (approximate). */
            size_t local = di->maxWorkGroup < 256 ? di->maxWorkGroup : 256;
            size_t lanes = 1u << 20;
            cl_uint perLane = (cl_uint)((uint64_t)words / 4ull / (uint64_t)lanes);
            double ms, bytes = (double)perLane * (double)lanes * 16.0;
            if (perLane == 0) { perLane = 1; lanes = (size_t)words / 4u; bytes = (double)lanes * 16.0; }
            CL_CHECK(clSetKernelArg(kStream, 0, sizeof(cl_mem), &dDs));
            CL_CHECK(clSetKernelArg(kStream, 1, sizeof(cl_uint), &perLane));
            CL_CHECK(clSetKernelArg(kStream, 2, sizeof(cl_mem), &dOut));
            ms = probeLaunch(dv, o, kStream, lanes, local, 3, -1, 0u);
            printf("| stream | %d | %llu | %llu | %u | %.3f | %.1f GB/s coalesced | (%.0f MiB read once) |\n", mib, (unsigned long long)local, (unsigned long long)lanes, perLane, ms, bytes / (ms / 1000.0) / 1e9, bytes / 1048576.0);
            fflush(stdout);
        }
        clReleaseMemObject(dDs);
    }
    {
        size_t local = di->maxWorkGroup < 256 ? di->maxWorkGroup : 256;
        size_t lanes = 1u << 20;
        cl_uint seed = 0x2468aceu;
        double ms, ops;
        CL_CHECK(clSetKernelArg(kAlu, 0, sizeof(cl_uint), &ALU_STEPS));
        CL_CHECK(clSetKernelArg(kAlu, 1, sizeof(cl_uint), &seed));
        CL_CHECK(clSetKernelArg(kAlu, 2, sizeof(cl_mem), &dOut));
        ms = probeLaunch(dv, o, kAlu, lanes, local, 3, 1, seed);
        ops = (double)lanes * (double)ALU_STEPS * 5.0;   /* mul, add, rotate, xor, add per step */
        printf("| alu | 0 | %llu | %llu | %u | %.3f | %.1f G int ops/s | %.3f G steps/s per compute unit (approximate: 5 ops per step counted) |\n",
               (unsigned long long)local, (unsigned long long)lanes, ALU_STEPS, ms, ops / (ms / 1000.0) / 1e9,
               (double)lanes * (double)ALU_STEPS / (ms / 1000.0) / 1e9 / (double)(di->computeUnits ? di->computeUnits : 1));
    }
    clReleaseMemObject(dOut);
    clReleaseKernel(kFill); clReleaseKernel(kChase); clReleaseKernel(kIndep); clReleaseKernel(kAlu); clReleaseKernel(kLine); clReleaseKernel(kStream); clReleaseKernel(kLine16);
    clReleaseProgram(prog);
    printf("memprobe: done\n");
    return 0;
}

int main(int argc, char** argv) {
    Options o = parseArgs(argc, argv);
    DeviceInfo* devs = NULL;
    int nDev, i, chosen = -1;
    DeviceInfo* di;
    Device dv;
    cl_int err = 0;
    size_t srcLen = 0;
    char* src;
    uint32_t nonces;
    cl_mem dOut;
    int sizes[5], nSizes = 0;
    SizeResult results[5];
    int cachePass = 1, overall, anyVec = 0;

#ifdef IGNEUM_CL_DYNAMIC
    if (!ig_cl_load()) { printf("%s %s\n", o.serve ? "error 0" : "FAIL:", ig_cl_error); fflush(stdout); return 2; }
#endif
    if (o.packDir) {
        /* Generic serve mode: the pack directory replaces the compiled-in pack (4 October 2026) */
        static char boundPath[1200];
        char perr[512];
        size_t n = strlen(o.packDir);
        if (!o.serve && !o.benchPack) { printf("FAIL: --pack goes with --serve or --bench-pack (the plain bench runs the compiled-in pack)\n"); return 2; }
        if (n > 1 && (o.packDir[n - 1] == '/' || o.packDir[n - 1] == '\\')) ((char*)o.packDir)[n - 1] = 0;
        if (!pf_load(o.packDir, &gPack, perr, sizeof(perr))) { printf("error 0 pack %s: %s\n", o.packDir, perr); fflush(stdout); return 2; }
        gGeneric = 1;
        gServeWords = 1u << gPack.datasetLog2;
#if IGNEUM_DATASET_MODE == 1
        gServeCacheWords = 1u << gPack.cacheLog2Words;
        gServeSegments = gPack.cacheSegments;
#endif
        if (!o.kernelGiven) { snprintf(boundPath, sizeof(boundPath), "%s/kernel_bound.cl", o.packDir); o.kernelPath = boundPath; o.kernelGiven = 1; }
    }
    printf("igneum-bench-cl  pack \"%s\"  (test harness: no pool, no network, no wallet)%s\n", gGeneric ? gPack.seedString : IGNEUM_SEED_STRING, gGeneric ? "  [--pack: generic serve mode]" : "");
    nDev = enumerateDevices(&devs);
    if (nDev == 0) { printf("%s no OpenCL platform or device found (is the GPU driver installed? it provides OpenCL)\n", o.serve ? "error 0" : "FAIL:"); fflush(stdout); return 2; }
    if (o.device >= 0) {
        if (o.device >= nDev) { printf("FAIL: --device %d out of range (%d devices)\n", o.device, nDev); return 2; }
        chosen = o.device;
    } else if (o.vendor) {
        for (i = 0; i < nDev; ++i) if ((devs[i].type & CL_DEVICE_TYPE_GPU) && devs[i].dupOf < 0 && strstr(devs[i].vendor, o.vendor)) { chosen = i; break; }
        if (chosen < 0) {
            printf("OpenCL devices (%d):\n", nDev);
            for (i = 0; i < nDev; ++i) printDevice(i, &devs[i], 0);
            printf("FAIL: no GPU whose vendor contains \"%s\"\n", o.vendor);
            return 2;
        }
    } else {
        for (i = 0; i < nDev; ++i) if ((devs[i].type & CL_DEVICE_TYPE_GPU) && devs[i].dupOf < 0) { chosen = i; break; }
        if (chosen < 0) chosen = 0;
    }
    printf("OpenCL devices (%d):\n", nDev);
    for (i = 0; i < nDev; ++i) printDevice(i, &devs[i], i == chosen && !o.list);
    {
        int dups = 0;
        for (i = 0; i < nDev; ++i) if (devs[i].dupOf >= 0) ++dups;
        if (dups) printf("platforms: %d device(s) hidden as the same card on an older platform of the same vendor (an old driver's OpenCL registration is still present)\n", dups);
    }
    if (o.list) return 0;
    di = &devs[chosen];
    printf("using device [%d] %s\n", chosen, di->name);
    if (di->dupOf >= 0) printf("NOTE: --device %d is the older platform's listing of the card [%d] (driver %s); results on it are for comparison only\n", chosen, di->dupOf, di->driver);
    if (o.timeWall < 0) o.timeWall = (strcmp(di->platformName, "Apple") == 0) ? 1 : 0;
    if (o.timeWall) printf("timing: host wall time (Apple's OpenCL event timestamps are not usable; the rate is still a device rate, see README.md)\n");
    else printf("timing: device event profiling (CL_PROFILING_COMMAND_START/END), like cudaEvent elapsed time\n");

    memset(&dv, 0, sizeof(dv));
    dv.ctx = clCreateContext(NULL, 1, &di->device, NULL, NULL, &err);
    CL_CHECK_ERR(err, "clCreateContext");
    dv.q = clCreateCommandQueue(dv.ctx, di->device, CL_QUEUE_PROFILING_ENABLE, &err);
    CL_CHECK_ERR(err, "clCreateCommandQueue");

    if (o.memprobe) {
        int rc = runMemprobe(&dv, di, &o);
        clReleaseCommandQueue(dv.q);
        clReleaseContext(dv.ctx);
        return rc;
    }
    if (o.benchPack && !o.packDir) { printf("FAIL: --bench-pack needs --pack <dir>\n"); return 2; }
    if (o.serve && !o.kernelGiven) {
        /* The bound kernel lives next to the compiled-in kernel.cl as kernel_bound.cl (packs from igneum-pow or igneum-miner export-pack). */
        static char boundPath[1024];
        size_t n = strlen(o.kernelPath);
        if (n >= 9 && strcmp(o.kernelPath + n - 9, "kernel.cl") == 0) {
            snprintf(boundPath, sizeof(boundPath), "%.*skernel_bound.cl", (int)(n - 9), o.kernelPath);
            o.kernelPath = boundPath;
        }
    }
    src = readFile(o.kernelPath, &srcLen);
    if (!src) { printf("FAIL: cannot read kernel source %s (run from proto-opencl/ or pass --kernel)\n", o.kernelPath); return 2; }
    printf("kernel source: %s (%llu bytes)\n", o.kernelPath, (unsigned long long)srcLen);
    setupProgram(&dv, di, &o, src, srcLen);
    free(src);
    printf("build options: %s\n", dv.buildOptions);
    printf("exchange: %s\n", dv.exchangeNote);
    if (o.serve) return runServe(&dv, di, &o);
    if (o.benchPack) return runBenchPack(&dv, di, &o);
    printKernelInfo(di, dv.kHash, "igneum_hash", dv.groupSize, "");
    printf("program: %d instructions x %d iterations, loads/hash %d, op mix %s\n",
           IGNEUM_INSTR_COUNT, IGNEUM_ITERATIONS, IGNEUM_LOADS_PER_HASH, IGNEUM_OP_MIX);
    printf("seed words: %08x %08x %08x %08x %08x %08x %08x %08x\n",
           SEEDW[0], SEEDW[1], SEEDW[2], SEEDW[3], SEEDW[4], SEEDW[5], SEEDW[6], SEEDW[7]);
    printf("day \"%s\" (d0 0x%08x, d1 0x%08x), pack dataset 2^%d words = %d MiB\n",
           IGNEUM_DAY_STRING, IGNEUM_DAY0, IGNEUM_DAY1, IGNEUM_DATASET_LOG2, packMib());
#if IGNEUM_DATASET_MODE == 1
#ifndef IGNEUM_MIXER_MULT
#define IGNEUM_MIXER_MULT 1
#endif
    printf("dataset construction: memory-hard (%u MiB ChaCha cache, %d dependent cache reads per 64-byte item, mixer x%d; proto-metal/MEMHARD.md)\n",
           (unsigned)(((uint64_t)CACHE_WORDS_HOST * 4u) >> 20), IGNEUM_ITEM_ROUNDS, IGNEUM_MIXER_MULT);
    cachePass = setupCache(&dv, di);
#else
    printf("dataset construction: closed-form ds_elem (the original prototype dataset, not memory-hard)\n");
#endif

    nonces = 1u << o.batchLog2;
    if (nonces % (32u * (uint32_t)o.groupWarps) != 0u) {
        printf("FAIL: 2^%d nonces is not a multiple of %d items per work-group\n", o.batchLog2, 32 * o.groupWarps);
        return 2;
    }
    dOut = clCreateBuffer(dv.ctx, CL_MEM_READ_WRITE, (size_t)nonces * sizeof(uint64_t), NULL, &err);
    CL_CHECK_ERR(err, "clCreateBuffer out");

    if (o.sweep) { sizes[0] = 4; sizes[1] = 64; sizes[2] = 256; sizes[3] = 512; sizes[4] = 1024; nSizes = 5; }
    else { sizes[0] = o.datasetMib; nSizes = 1; }
    for (i = 0; i < nSizes; ++i) results[i] = runSize(&dv, di, &o, sizes[i], dOut, nonces);
    CL_CHECK(clReleaseMemObject(dOut));

    printf("\n=== summary (%s, %s, pack %s, batch 2^%d x %d, %d unit(s)/work-group, exchange %s, %s time) ===\n",
           di->name, di->platformName, IGNEUM_SEED_STRING, o.batchLog2, o.batches, o.groupWarps, exchangeName(dv.exchange), o.timeWall ? "wall" : "device-event");
    printf("| dataset MiB | %s ms (second) | Mhash/s | GB/s useful | random loads/s (G) | loads/hash | dataset self-test | vectors |\n",
           IGNEUM_DATASET_MODE == 1 ? "build" : "fill");
    printf("|---|---|---|---|---|---|---|---|\n");
    overall = cachePass;
    for (i = 0; i < nSizes; ++i) {
        const SizeResult* r = &results[i];
        overall = overall && r->dsPass && (!r->vecChecked || r->vecPass);
        anyVec = anyVec || r->vecChecked;
        printf("| %d | %.2f | %.3f | %.2f | %.2f | %d | %s | %s |\n",
               r->mib, r->fillSecondMs, r->hashesPerSec / 1e6, r->gbps,
               r->hashesPerSec * (double)IGNEUM_LOADS_PER_HASH / 1e9, IGNEUM_LOADS_PER_HASH,
               r->dsPass ? "PASS" : "FAIL",
               r->vecChecked ? (r->vecPass ? "PASS (3 warps, standalone and in batch)" : "FAIL") : "skipped (not pack size)");
    }
#if IGNEUM_DATASET_MODE == 1
    printf("cache: device fill %.2f ms (second), host fill %.1f ms one thread, cache check %s\n", gCacheFillSecondMs, gCacheHostMs, cachePass ? "PASS" : "FAIL");
    if (gCache) clReleaseMemObject(gCache);
    free(hCache);
#endif
    if (gProfilingFailures) printf("NOTE: %d event profiling queries failed on this runtime (timings printed as -1.00 ms are unavailable)\n", gProfilingFailures);
    printf("exchange: %s\n", dv.exchangeNote);
    if (!anyVec) printf("NOTE: no vectors were checked. Run at %d MiB (the default) to verify against the Mac.\n", packMib());
    printf("OVERALL: %s\n", overall ? "PASS" : "FAIL");

    releaseProgram(&dv);
    clReleaseCommandQueue(dv.q);
    clReleaseContext(dv.ctx);
    for (i = 0; i < nDev; ++i) free(devs[i].extensions);
    free(devs);
    return overall ? 0 : 1;
}