/*
 * keccak_miner.cu - Generic CUDA Keccak-256 PoW miner, v3 (zero-loss).
 *
 * Generic = any prefix length, decided at runtime, with NO hashrate
 * loss vs a fixed-length build. The trick (validated by production
 * miners, e.g. Kaspa kHeavyHash, ERADICATE2, nft-mint-rarity-toolkit):
 * everything constant for a job is computed ONCE on the host and folded
 * into the state; the kernel's hot loop only touches the 1-2 lanes the
 * nonce actually changes.
 *
 * Preimage: keccak256(prefix || nonce_be32), nonce = uint64 in the low
 * 8 bytes of a 32-byte big-endian field. Prefix length is arbitrary
 * (0..1024 bytes). Examples:
 *   AMNT: prefix = sessionSeedBase(32) || wallet(32) || agentId(32)
 *           || lane(32) || day(32)                      = 160 bytes
 *   GEM:  prefix = seed(32) || wallet(32, address left-padded) = 64 bytes
 *
 * Per job (host, once), prepare_params() computes:
 *   - mid: Keccak state after absorbing all complete 136-byte blocks
 *     of the prefix, with the FIRST final block's constant lanes
 *     (prefix remainder + zero nonce field + padding, if single-block)
 *     already XORed into lanes 0..16.
 *   - c1: second final block's constant lanes (two-block case only).
 *   - lane/shift/spill descriptors telling the kernel exactly which
 *     1-2 lanes the nonce uint64 touches and at what bit shift.
 * Per hash (kernel): load mid, XOR the nonce into 1-2 lanes
 * (bswap64 + shift), run 1 permutation (2 in the two-block case),
 * compare to target. No per-hash lane rebuilding, no length branch
 * in the hot loop (two_blocks / lane / shift are uniform per launch).
 *
 * Dynamic indexing pitfall: lane numbers are runtime values, so the
 * kernel applies the nonce XOR through xor_lane_pair(), a switch with
 * static indices per case. Indexing state[] by a runtime value would
 * spill the 25-lane state to local memory and kill performance.
 *
 * Correctness: the Keccak code is shared with the CPU test build
 * (same __host__ __device__ functions). Verify with:
 *   gcc -O2 -DKECCAK_MINER_HOST_TEST -x c -o keccak_miner_test keccak_miner.cu
 *   ./keccak_miner_test vectors
 *   ./keccak_miner_test selftest
 * selftest checks split==full for EVERY prefix length 0..271 plus
 * 1000 (covers all remainder values 0..135 and both sides of the
 * one-block/two-block boundary) x 8 nonces.
 *
 * Compile (needs CUDA toolkit):
 *   nvcc -O3 -gencode arch=compute_89,code=sm_89 \
 *        -gencode arch=compute_120,code=sm_120 \
 *        -gencode arch=compute_120,code=compute_120 \
 *        -cudart static -o keccak_miner keccak_miner.cu
 * Target machine needs only an NVIDIA driver (no CUDA toolkit).
 *
 * Usage:
 *   ./keccak_miner <prefix_hex> <target_hex_32B> <count> [start_nonce]
 *   prefix_hex: any even length, 0..1024 bytes (0x prefix optional).
 *   Output: JSON {"nonces":[...],"hashes":[...]} on stdout, stats on stderr.
 */
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <stdint.h>

#ifdef __CUDACC__
#include <cuda_runtime.h>
#define HD __host__ __device__
#else
#define HD
#endif

#define MAX_PREFIX_BYTES 1024
#define KECCAK_RATE 136

/* ------------------------------------------------------------------ */
/* Keccak-f[1600], optimized.                                         */
/* ------------------------------------------------------------------ */

HD static inline uint64_t rol64(uint64_t x, int n) {
    return n ? (x << n) | (x >> (64 - n)) : x;
}

/*
 * One Keccak-f1600 permutation. State is 25 uint64_t lanes, index = x + 5*y.
 * Rho+Pi is done with a 5-lane rotating buffer (no 25-lane temp copy).
 */
HD static void keccak_f1600_opt(uint64_t st[25]) {
    static const uint64_t RC[24] = {
        0x0000000000000001ULL, 0x0000000000008082ULL, 0x800000000000808aULL,
        0x8000000080008000ULL, 0x000000000000808bULL, 0x0000000080000001ULL,
        0x8000000080008081ULL, 0x8000000000008009ULL, 0x000000000000008aULL,
        0x0000000000000088ULL, 0x0000000080008009ULL, 0x000000008000000aULL,
        0x000000008000808bULL, 0x800000000000008bULL, 0x8000000000008089ULL,
        0x8000000000008003ULL, 0x8000000000008002ULL, 0x8000000000000080ULL,
        0x000000000000800aULL, 0x800000008000000aULL, 0x8000000080008081ULL,
        0x8000000000008080ULL, 0x0000000080000001ULL, 0x8000000080008008ULL
    };
    /* Rotation offsets r[x][y]. */
    static const int ROT[5][5] = {
        {  0, 36,  3, 41, 18 },
        {  1, 44, 10, 45,  2 },
        { 62,  6, 43, 15, 61 },
        { 28, 55, 25, 21, 56 },
        { 27, 20, 39,  8, 14 }
    };

    #pragma unroll
    for (int round = 0; round < 24; round++) {
        uint64_t C[5], D[5];
        #pragma unroll
        for (int x = 0; x < 5; x++)
            C[x] = st[x] ^ st[x + 5] ^ st[x + 10] ^ st[x + 15] ^ st[x + 20];
        #pragma unroll
        for (int x = 0; x < 5; x++)
            D[x] = C[(x + 4) % 5] ^ rol64(C[(x + 1) % 5], 1);
        #pragma unroll
        for (int x = 0; x < 5; x++)
            #pragma unroll
            for (int y = 0; y < 5; y++)
                st[x + 5 * y] ^= D[x];

        /* Rho + Pi. B is fully unrolled and consumed immediately, so the
           compiler keeps it in registers (no local-memory spill). */
        uint64_t B[25];
        #pragma unroll
        for (int x = 0; x < 5; x++)
            #pragma unroll
            for (int y = 0; y < 5; y++)
                B[y + 5 * ((2 * x + 3 * y) % 5)] = rol64(st[x + 5 * y], ROT[x][y]);

        /* Chi. */
        #pragma unroll
        for (int x = 0; x < 5; x++)
            #pragma unroll
            for (int y = 0; y < 5; y++)
                st[x + 5 * y] = B[x + 5 * y] ^ ((~B[((x + 1) % 5) + 5 * y]) & B[((x + 2) % 5) + 5 * y]);

        /* Iota. */
        st[0] ^= RC[round];
    }
}

/* Keccak-256: rate 136 bytes, suffix 0x01, final byte OR 0x80. */
HD static void keccak256(const uint8_t *in, size_t inlen, uint8_t out[32]) {
    uint64_t st[25];
    memset(st, 0, sizeof(st));
    const size_t rate = KECCAK_RATE;

    while (inlen >= rate) {
        for (size_t i = 0; i < rate / 8; i++) {
            uint64_t lane = 0;
            for (int b = 0; b < 8; b++)
                lane |= (uint64_t)in[i * 8 + b] << (8 * b);
            st[i] ^= lane;
        }
        keccak_f1600_opt(st);
        in += rate;
        inlen -= rate;
    }

    uint8_t block[KECCAK_RATE];
    memset(block, 0, sizeof(block));
    memcpy(block, in, inlen);
    block[inlen] ^= 0x01;
    block[rate - 1] ^= 0x80;
    for (size_t i = 0; i < rate / 8; i++) {
        uint64_t lane = 0;
        for (int b = 0; b < 8; b++)
            lane |= (uint64_t)block[i * 8 + b] << (8 * b);
        st[i] ^= lane;
    }
    keccak_f1600_opt(st);

    for (int i = 0; i < 4; i++)
        for (int b = 0; b < 8; b++)
            out[i * 8 + b] = (uint8_t)(st[i] >> (8 * b));
}

/* True if hash <= target. Both are 32-byte big-endian. Compare as 4x u64. */
HD static int hash_le_target(const uint8_t *hash, const uint8_t *target) {
    /* Big-endian: first byte is most significant. Load as u64 big-endian. */
    #pragma unroll
    for (int w = 0; w < 4; w++) {
        uint64_t h = 0, t = 0;
        #pragma unroll
        for (int b = 0; b < 8; b++) {
            h = (h << 8) | hash[w * 8 + b];
            t = (t << 8) | target[w * 8 + b];
        }
        if (h < t) return 1;
        if (h > t) return 0;
    }
    return 1;
}

/*
 * Absorb all complete 136-byte blocks of prefix into mid; the leftover
 * bytes (0..135) go to tail. Works for any prefix length, including 0.
 */
HD static void absorb_prefix(const uint8_t *prefix, size_t plen,
                             uint64_t mid[25], uint8_t *tail, int *tail_len) {
    memset(mid, 0, 25 * sizeof(uint64_t));
    size_t off = 0;
    while (plen - off >= (size_t)KECCAK_RATE) {
        for (int i = 0; i < 17; i++) {
            uint64_t lane = 0;
            for (int b = 0; b < 8; b++)
                lane |= (uint64_t)prefix[off + (size_t)i * 8 + (size_t)b] << (8 * b);
            mid[i] ^= lane;
        }
        keccak_f1600_opt(mid);
        off += (size_t)KECCAK_RATE;
    }
    *tail_len = (int)(plen - off);
    if (*tail_len > 0) memcpy(tail, prefix + off, (size_t)*tail_len);
}

/*
 * Build the 17 rate lanes of one block from data_len data bytes
 * (<= 136), then if pad: byte data_len ^= 0x01, byte 135 ^= 0x80
 * (they merge to 0x81 when data_len == 135). Caller guarantees
 * data_len < 136 when pad is set.
 */
HD static void build_block_lanes_const(const uint8_t *data, int data_len,
                                       int pad, uint64_t blk[17]) {
    uint8_t block[KECCAK_RATE];
    memset(block, 0, sizeof(block));
    if (data_len > 0) memcpy(block, data, (size_t)data_len);
    if (pad) {
        block[data_len] ^= 0x01;
        block[KECCAK_RATE - 1] ^= 0x80;
    }
    for (int i = 0; i < 17; i++) {
        uint64_t lane = 0;
        for (int b = 0; b < 8; b++)
            lane |= (uint64_t)block[i * 8 + b] << (8 * b);
        blk[i] = lane;
    }
}

typedef struct {
    /* Mid-state with block-1 constant lanes folded into lanes 0..16. */
    uint64_t mid[25];
    /* Block-2 constant lanes (two_blocks only). */
    uint64_t c1[17];
    int      two_blocks; /* 1 when tail_len + 32 >= 136 */
    /* Block-1 nonce placement: uint64 at stream bytes [R+24, R+32).
       lane1 < 0 means no nonce byte in block 1. spill1: the nonce
       also touches lane1+1 (only when fully inside block 1). */
    int      lane1;
    int      shift1; /* left shift bits for bswap64(nonce) */
    int      spill1;
    /* Block-2 nonce placement. kind2: 0 none, 1 left-shift form
       (nonce fully in block 2), 2 right-shift form (straddle:
       st[0] ^= bswap64(nonce) >> shift2). */
    int      lane2;
    int      shift2;
    int      kind2;
    int      spill2; /* kind2 == 1 only */
    uint8_t  target[32];
} mine_params_t;

/*
 * Per-job setup (host, once). Fills params from prefix + target.
 * The nonce uint64 sits at stream bytes [R+24, R+32) where R is the
 * prefix remainder length; its big-endian bytes map to lane bits via
 * w = bswap64(nonce): st[lane] ^= w << shift (spill to lane+1 when
 * the 8 bytes straddle a lane boundary inside the same block).
 */
HD static void prepare_params(const uint8_t *prefix, size_t plen,
                              const uint8_t *target, mine_params_t *p) {
    memset(p, 0, sizeof(*p));
    uint64_t mid[25];
    uint8_t tail[KECCAK_RATE];
    int R = 0;
    absorb_prefix(prefix, plen, mid, tail, &R);
    int total = R + 32;
    int two_blocks = (total >= KECCAK_RATE);
    p->two_blocks = two_blocks;

    if (!two_blocks) {
        /* Single final block: tail + 32 zero nonce bytes + padding. */
        uint8_t b1[KECCAK_RATE];
        memset(b1, 0, sizeof(b1));
        if (R > 0) memcpy(b1, tail, (size_t)R);
        uint64_t c0[17];
        build_block_lanes_const(b1, total, 1, c0);
        for (int i = 0; i < 17; i++) p->mid[i] = mid[i] ^ c0[i];
        for (int i = 17; i < 25; i++) p->mid[i] = mid[i];
    } else {
        /* Block 1: tail + first (136 - R) nonce bytes (zeros), no pad. */
        uint8_t b1[KECCAK_RATE];
        memset(b1, 0, sizeof(b1));
        if (R > 0) memcpy(b1, tail, (size_t)R);
        uint64_t c0[17];
        build_block_lanes_const(b1, KECCAK_RATE, 0, c0);
        for (int i = 0; i < 17; i++) p->mid[i] = mid[i] ^ c0[i];
        for (int i = 17; i < 25; i++) p->mid[i] = mid[i];
        /* Block 2: remaining nonce bytes (zeros) + padding. */
        int rem = total - KECCAK_RATE;
        uint8_t b2[KECCAK_RATE];
        memset(b2, 0, sizeof(b2));
        build_block_lanes_const(b2, rem, 1, p->c1);
    }

    /* Nonce placement. */
    p->lane1 = -1; p->shift1 = 0; p->spill1 = 0;
    p->lane2 = -1; p->shift2 = 0; p->kind2 = 0; p->spill2 = 0;
    if (R + 24 < KECCAK_RATE) {
        p->lane1 = (R + 24) / 8;
        p->shift1 = ((R + 24) % 8) * 8;
        p->spill1 = (R + 32 <= KECCAK_RATE && p->shift1 != 0) ? 1 : 0;
    }
    if (two_blocks) {
        if (R + 24 >= KECCAK_RATE) {
            /* Nonce fully in block 2. */
            int pos2 = R + 24 - KECCAK_RATE;
            p->lane2 = pos2 / 8;
            p->shift2 = (pos2 % 8) * 8;
            p->kind2 = 1;
            p->spill2 = (p->shift2 != 0) ? 1 : 0;
        } else if (R + 32 > KECCAK_RATE) {
            /* Straddle: low bytes of the nonce land in block 2 lane 0. */
            int k1 = KECCAK_RATE - R - 24; /* nonce bytes in block 1 */
            p->lane2 = 0;
            p->shift2 = k1 * 8; /* st[0] ^= bswap64(nonce) >> shift2 */
            p->kind2 = 2;
        }
        /* else (R + 32 == 136): nonce fully in block 1, block 2 is
           padding only; kind2 stays 0. */
    }
    memcpy(p->target, target, 32);
}

/* Byte-swap a 64-bit word. Written with shifts/ORs (not
   __builtin_bswap64, which nvcc rejects in __device__ code):
   the CPU test build verifies this exact code path, and the
   compiler folds it to one bswap/PRMT instruction. */
HD static inline uint64_t bswap64_u64(uint64_t x) {
    return ((x & 0x00000000000000ffULL) << 56) | ((x & 0x000000000000ff00ULL) << 40) |
           ((x & 0x0000000000ff0000ULL) << 24) | ((x & 0x00000000ff000000ULL) << 8) |
           ((x & 0x000000ff00000000ULL) >> 8)  | ((x & 0x0000ff0000000000ULL) >> 24) |
           ((x & 0x00ff000000000000ULL) >> 40) | ((x & 0xff00000000000000ULL) >> 56);
}

/*
 * XOR (w << shift) into st[lane] and (w >> (64 - shift)) into
 * st[lane+1]. Written as a switch with static indices per case:
 * indexing st[] by a runtime value would spill the 25-lane state
 * to local memory. hi must be 0 when there is no spill.
 * lane is guaranteed 0..16 here; lane+1 <= 17 is a valid index.
 */
HD static inline void xor_lane_pair(uint64_t st[25], int lane,
                                    uint64_t lo, uint64_t hi) {
    switch (lane) {
        case 0:  st[0]  ^= lo; st[1]  ^= hi; break;
        case 1:  st[1]  ^= lo; st[2]  ^= hi; break;
        case 2:  st[2]  ^= lo; st[3]  ^= hi; break;
        case 3:  st[3]  ^= lo; st[4]  ^= hi; break;
        case 4:  st[4]  ^= lo; st[5]  ^= hi; break;
        case 5:  st[5]  ^= lo; st[6]  ^= hi; break;
        case 6:  st[6]  ^= lo; st[7]  ^= hi; break;
        case 7:  st[7]  ^= lo; st[8]  ^= hi; break;
        case 8:  st[8]  ^= lo; st[9]  ^= hi; break;
        case 9:  st[9]  ^= lo; st[10] ^= hi; break;
        case 10: st[10] ^= lo; st[11] ^= hi; break;
        case 11: st[11] ^= lo; st[12] ^= hi; break;
        case 12: st[12] ^= lo; st[13] ^= hi; break;
        case 13: st[13] ^= lo; st[14] ^= hi; break;
        case 14: st[14] ^= lo; st[15] ^= hi; break;
        case 15: st[15] ^= lo; st[16] ^= hi; break;
        case 16: st[16] ^= lo; st[17] ^= hi; break;
        default: break;
    }
}

/*
 * Finish one preimage from params + nonce. Shared by the kernel and
 * the CPU test build, so the test verifies the exact hot-loop code.
 */
HD static inline void finish_nonce(const uint64_t *mid, const mine_params_t *p,
                                   uint64_t nonce, uint8_t out[32]) {
    uint64_t st[25];
    for (int i = 0; i < 25; i++) st[i] = mid[i];
    uint64_t w = bswap64_u64(nonce);
    if (p->lane1 >= 0) {
        uint64_t hi = p->spill1 ? (w >> (64 - p->shift1)) : 0;
        xor_lane_pair(st, p->lane1, w << p->shift1, hi);
    }
    keccak_f1600_opt(st);
    if (p->two_blocks) {
        #pragma unroll
        for (int i = 0; i < 17; i++) st[i] ^= p->c1[i];
        if (p->kind2 == 1) {
            uint64_t hi = p->spill2 ? (w >> (64 - p->shift2)) : 0;
            xor_lane_pair(st, p->lane2, w << p->shift2, hi);
        } else if (p->kind2 == 2) {
            st[0] ^= w >> p->shift2;
        }
        keccak_f1600_opt(st);
    }
    for (int i = 0; i < 4; i++)
        for (int b = 0; b < 8; b++)
            out[i * 8 + b] = (uint8_t)(st[i] >> (8 * b));
}

/* ------------------------------------------------------------------ */
/* Shared helpers (plain C).                                           */
/* ------------------------------------------------------------------ */

static int hexval(char c) {
    if (c >= '0' && c <= '9') return c - '0';
    if (c >= 'a' && c <= 'f') return c - 'a' + 10;
    if (c >= 'A' && c <= 'F') return c - 'A' + 10;
    return -1;
}

/* Hex string -> bytes. Returns byte count, or -1 on error. */
static int hex_to_bytes_any(const char *hex, uint8_t *out, size_t max) {
    if (hex[0] == '0' && (hex[1] == 'x' || hex[1] == 'X')) hex += 2;
    size_t len = strlen(hex);
    if (len % 2 != 0 || len / 2 > max) return -1;
    for (size_t i = 0; i < len / 2; i++) {
        int hi = hexval(hex[2 * i]), lo = hexval(hex[2 * i + 1]);
        if (hi < 0 || lo < 0) return -1;
        out[i] = (uint8_t)((hi << 4) | lo);
    }
    return (int)(len / 2);
}

static int hex_to_bytes(const char *hex, uint8_t *out, size_t want) {
    int n = hex_to_bytes_any(hex, out, want);
    return n == (int)want ? 0 : -1;
}

static void to_hex(const uint8_t *in, size_t len, char *out) {
    static const char hexd[] = "0123456789abcdef";
    for (size_t i = 0; i < len; i++) {
        out[2 * i] = hexd[in[i] >> 4];
        out[2 * i + 1] = hexd[in[i] & 0x0f];
    }
    out[2 * len] = '\0';
}

#ifdef __CUDACC__
/* ------------------------------------------------------------------ */
/* CUDA kernel + main. 4-way ILP, no per-hash atomic.                  */
/* ------------------------------------------------------------------ */

#define ILP 4
#define CHECK_EVERY 1024

__global__ void mine_kernel(const mine_params_t *params, uint64_t start_nonce,
                            unsigned long long *found_count,
                            uint64_t *found_nonces, uint8_t *found_hashes,
                            int max_count) {
    uint64_t tid = blockIdx.x * (uint64_t)blockDim.x + threadIdx.x;
    uint64_t stride = (uint64_t)gridDim.x * blockDim.x * ILP;

    /* Load mid-state into registers once; reuse across all hashes. */
    uint64_t mid[25];
    #pragma unroll
    for (int i = 0; i < 25; i++) mid[i] = params->mid[i];

    uint64_t base = start_nonce + tid * ILP;
    unsigned int since_check = 0;

    for (;;) {
        /* Interleave ILP independent hashes. */
        uint8_t hashes[ILP][32];
        uint64_t nonces[ILP];
        #pragma unroll
        for (int k = 0; k < ILP; k++) {
            nonces[k] = base + (uint64_t)k * (uint64_t)gridDim.x * blockDim.x;
            finish_nonce(mid, params, nonces[k], hashes[k]);
        }

        #pragma unroll
        for (int k = 0; k < ILP; k++) {
            if (hash_le_target(hashes[k], params->target)) {
                unsigned long long idx = atomicAdd(found_count, 1ULL);
                if (idx < (unsigned long long)max_count) {
                    found_nonces[idx] = nonces[k];
                    for (int i = 0; i < 32; i++) found_hashes[idx * 32 + i] = hashes[k][i];
                }
                if (idx + 1 >= (unsigned long long)max_count) return;
            }
        }

        base += stride;
        /* Cheap exit check: only touch global memory once per 1024 rounds. */
        if (++since_check >= CHECK_EVERY) {
            since_check = 0;
            if (*found_count >= (unsigned long long)max_count) return;
        }
    }
}

#define CUDA_CHECK(call) do { \
    cudaError_t _e = (call); \
    if (_e != cudaSuccess) { \
        fprintf(stderr, "CUDA error at %s:%d: %s\n", __FILE__, __LINE__, cudaGetErrorString(_e)); \
        return 1; \
    } \
} while (0)

int main(int argc, char **argv) {
    if (argc < 4 || argc > 5) {
        fprintf(stderr,
            "Usage: %s <prefix_hex> <target_hex_32B> <count> [start_nonce]\n"
            "  prefix_hex: any even length, 0..%d bytes (0x prefix optional)\n",
            argv[0], MAX_PREFIX_BYTES);
        return 2;
    }
    uint8_t prefix[MAX_PREFIX_BYTES], target[32];
    int plen = hex_to_bytes_any(argv[1], prefix, MAX_PREFIX_BYTES);
    if (plen < 0) { fprintf(stderr, "bad prefix hex (even length, <= %d bytes)\n", MAX_PREFIX_BYTES); return 2; }
    if (hex_to_bytes(argv[2], target, 32) != 0) { fprintf(stderr, "bad target hex (need 32 bytes)\n"); return 2; }
    int count = atoi(argv[3]);
    if (count <= 0 || count > 1024) { fprintf(stderr, "count 1..1024\n"); return 2; }
    uint64_t start_nonce = (argc >= 5) ? strtoull(argv[4], NULL, 10) : 0;

    mine_params_t h_params;
    prepare_params(prefix, (size_t)plen, target, &h_params);

    mine_params_t *d_params = NULL;
    unsigned long long *d_count = NULL;
    uint64_t *d_nonces = NULL;
    uint8_t *d_hashes = NULL;
    CUDA_CHECK(cudaMalloc((void **)&d_params, sizeof(mine_params_t)));
    CUDA_CHECK(cudaMemcpy(d_params, &h_params, sizeof(mine_params_t), cudaMemcpyHostToDevice));
    CUDA_CHECK(cudaMalloc((void **)&d_count, sizeof(unsigned long long)));
    CUDA_CHECK(cudaMemset(d_count, 0, sizeof(unsigned long long)));
    CUDA_CHECK(cudaMalloc((void **)&d_nonces, sizeof(uint64_t) * (size_t)count));
    CUDA_CHECK(cudaMalloc((void **)&d_hashes, (size_t)32 * (size_t)count));

    int dev = 0;
    CUDA_CHECK(cudaGetDevice(&dev));
    cudaDeviceProp prop;
    CUDA_CHECK(cudaGetDeviceProperties(&prop, dev));
    int block = 256;
    int grid = prop.multiProcessorCount > 0 ? prop.multiProcessorCount * 8 : 1024;
    fprintf(stderr, "device: %s, SMs=%d, block=%d, grid=%d, threads=%lld (ILP=%d), prefix=%dB two_blocks=%d lane1=%d shift1=%d\n",
            prop.name, prop.multiProcessorCount, block, grid,
            (long long)block * grid, ILP, plen, h_params.two_blocks,
            h_params.lane1, h_params.shift1);

    cudaEvent_t t0, t1;
    CUDA_CHECK(cudaEventCreate(&t0));
    CUDA_CHECK(cudaEventCreate(&t1));
    CUDA_CHECK(cudaEventRecord(t0));
    mine_kernel<<<grid, block>>>(d_params, start_nonce, d_count, d_nonces, d_hashes, count);
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaDeviceSynchronize());
    CUDA_CHECK(cudaEventRecord(t1));
    CUDA_CHECK(cudaEventSynchronize(t1));
    float ms = 0.0f;
    CUDA_CHECK(cudaEventElapsedTime(&ms, t0, t1));

    unsigned long long h_count = 0;
    CUDA_CHECK(cudaMemcpy(&h_count, d_count, sizeof(unsigned long long), cudaMemcpyDeviceToHost));
    if (h_count > (unsigned long long)count) h_count = (unsigned long long)count;
    uint64_t *h_nonces = (uint64_t *)malloc(sizeof(uint64_t) * (size_t)count);
    uint8_t *h_hashes = (uint8_t *)malloc((size_t)32 * (size_t)count);
    if (!h_nonces || !h_hashes) { fprintf(stderr, "host malloc failed\n"); return 1; }
    CUDA_CHECK(cudaMemcpy(h_nonces, d_nonces, sizeof(uint64_t) * (size_t)h_count, cudaMemcpyDeviceToHost));
    CUDA_CHECK(cudaMemcpy(h_hashes, d_hashes, (size_t)32 * (size_t)h_count, cudaMemcpyDeviceToHost));

    printf("{\"nonces\":[");
    for (unsigned long long i = 0; i < h_count; i++)
        printf("%s\"%llu\"", i ? "," : "", (unsigned long long)h_nonces[i]);
    printf("],\"hashes\":[");
    for (unsigned long long i = 0; i < h_count; i++) {
        char hex[65];
        to_hex(h_hashes + i * 32, 32, hex);
        printf("%s\"0x%s\"", i ? "," : "", hex);
    }
    printf("]}\n");
    fprintf(stderr, "found=%llu/%d elapsed=%.3fs\n", h_count, count, ms / 1000.0);

    free(h_nonces);
    free(h_hashes);
    cudaFree(d_params);
    cudaFree(d_count);
    cudaFree(d_nonces);
    cudaFree(d_hashes);
    return h_count == (unsigned long long)count ? 0 : 1;
}

#elif defined(KECCAK_MINER_HOST_TEST)
/* CPU test build: verifies the exact code the kernel uses, as host. */

static int test_vector(const char *input, const char *expected_hex) {
    uint8_t hash[32];
    char hex[65];
    keccak256((const uint8_t *)input, strlen(input), hash);
    to_hex(hash, 32, hex);
    int ok = strcmp(hex, expected_hex) == 0;
    printf("vector keccak256(\"%s\") = 0x%s  %s\n", input, hex, ok ? "PASS" : "FAIL");
    return ok;
}

static void full_hash(const uint8_t *prefix, size_t plen, uint64_t nonce, uint8_t out[32]) {
    uint8_t *buf = (uint8_t *)malloc(plen + 32);
    memcpy(buf, prefix, plen);
    memset(buf + plen, 0, 32);
    uint64_t v = nonce;
    for (int i = 31; i >= 0 && v; i--) { buf[plen + (size_t)i] = (uint8_t)(v & 0xff); v >>= 8; }
    keccak256(buf, plen + 32, out);
    free(buf);
}

static int test_split_vs_full(const uint8_t *prefix, size_t plen, uint64_t nonce) {
    mine_params_t p;
    uint8_t target[32];
    memset(target, 0xff, 32);
    prepare_params(prefix, plen, target, &p);
    uint8_t h_split[32], h_full[32];
    finish_nonce(p.mid, &p, nonce, h_split);
    full_hash(prefix, plen, nonce, h_full);
    int ok = memcmp(h_split, h_full, 32) == 0;
    if (!ok) {
        char hs[65], hf[65];
        to_hex(h_split, 32, hs);
        to_hex(h_full, 32, hf);
        printf("split==full plen=%zu nonce=%llu  FAIL\n  split 0x%s\n  full  0x%s\n",
               plen, (unsigned long long)nonce, hs, hf);
    }
    return ok;
}

int main(int argc, char **argv) {
    if (argc < 2) {
        fprintf(stderr,
            "Usage:\n"
            "  %s vectors\n"
            "  %s hash <prefix_hex> <nonce>\n"
            "  %s split <prefix_hex> <nonce>\n"
            "  %s selftest\n", argv[0], argv[0], argv[0], argv[0]);
        return 2;
    }

    if (strcmp(argv[1], "vectors") == 0) {
        int ok = 1;
        ok &= test_vector("", "c5d2460186f7233c927e7db2dcc703c0e500b653ca82273b7bfad8045d85a470");
        ok &= test_vector("abc", "4e03657aea45a94fc7d47ba826c8d667c0d1e6e33a64a036ec44f58fa12d6c45");
        printf("vectors: %s\n", ok ? "ALL PASS" : "FAIL");
        return ok ? 0 : 1;
    }

    if (strcmp(argv[1], "selftest") == 0) {
        int ok = 1;
        ok &= test_vector("", "c5d2460186f7233c927e7db2dcc703c0e500b653ca82273b7bfad8045d85a470");
        ok &= test_vector("abc", "4e03657aea45a94fc7d47ba826c8d667c0d1e6e33a64a036ec44f58fa12d6c45");
        /* Every prefix length 0..271 (all remainders 0..135, both sides
           of the one/two-block boundary at R=104) plus 1000. */
        static const uint64_t nonces[] = {0, 1, 2, 3, 7, 12345, 999999,
                                          4294967296ULL, 1099511627776ULL,
                                          9007199254740991ULL};
        uint8_t prefix[MAX_PREFIX_BYTES];
        for (size_t i = 0; i < MAX_PREFIX_BYTES; i++)
            prefix[i] = (uint8_t)((i * 37 + 11) & 0xff);
        int total = 0, passed = 0;
        for (size_t plen = 0; plen <= 271; plen++) {
            for (size_t n = 0; n < sizeof(nonces) / sizeof(nonces[0]); n++) {
                total++;
                if (test_split_vs_full(prefix, plen, nonces[n])) passed++;
                else ok = 0;
            }
        }
        for (size_t n = 0; n < sizeof(nonces) / sizeof(nonces[0]); n++) {
            total++;
            if (test_split_vs_full(prefix, 1000, nonces[n])) passed++;
            else ok = 0;
        }
        printf("selftest: %d/%d split==full  %s\n", passed, total,
               ok ? "ALL PASS" : "FAIL");
        return ok ? 0 : 1;
    }

    if (argc < 4) {
        fprintf(stderr, "need <prefix_hex> <nonce>\n");
        return 2;
    }
    uint8_t prefix[MAX_PREFIX_BYTES];
    int plen = hex_to_bytes_any(argv[2], prefix, MAX_PREFIX_BYTES);
    if (plen < 0) { fprintf(stderr, "bad prefix hex\n"); return 2; }
    uint64_t nonce = strtoull(argv[3], NULL, 10);

    if (strcmp(argv[1], "hash") == 0) {
        uint8_t h[32];
        char hex[65];
        full_hash(prefix, (size_t)plen, nonce, h);
        to_hex(h, 32, hex);
        printf("0x%s\n", hex);
        return 0;
    }
    if (strcmp(argv[1], "split") == 0) {
        mine_params_t p;
        uint8_t target[32];
        memset(target, 0xff, 32);
        prepare_params(prefix, (size_t)plen, target, &p);
        uint8_t h[32];
        char hex[65];
        finish_nonce(p.mid, &p, nonce, h);
        to_hex(h, 32, hex);
        printf("0x%s\n", hex);
        return 0;
    }

    fprintf(stderr, "unknown mode: %s\n", argv[1]);
    return 2;
}

#else
#error "Compile with nvcc for the CUDA miner, or with gcc -DKECCAK_MINER_HOST_TEST -x c for the CPU test build."
#endif