/*
 * Bitcoin P2PK Bruteforce - VERSIONE GPU (CUDA)
 *
 * Stessa strategia della versione CPU (p2pk_bruteforce.cpp): 1 moltiplicazione
 * scalare "costosa" seguita da un batch di addizioni EC economiche in
 * coordinate Jacobiane con Z1=1, poi UNA sola inversione di campo (trucco di
 * Montgomery) per convertire l'intero batch in affine. La differenza è che
 * qui non c'è GMP né libsecp256k1 nel loop caldo: l'aritmetica di campo
 * mod p (256 bit) è implementata a mano in CUDA (limb a 64 bit + __int128),
 * perché GMP non gira su device, e viene eseguita in parallelo da decine di
 * migliaia di thread GPU invece che da una manciata di thread CPU.
 *
 * DISCLAIMER: Solo per scopi educativi e di ricerca.
 */

#include <cuda_runtime.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <stdint.h>
#include <time.h>
#include <signal.h>
#include <secp256k1.h>
#include <vector>
#include <array>
#include <algorithm>
#include <fstream>
#include <string>
#include <cctype>

// ============================================================================
// CONFIGURAZIONE
// ============================================================================

#define GPU_EC_BATCH_SIZE 128           // Chiavi generate per moltiplicazione scalare
#define GPU_EC_BATCH_MULT (GPU_EC_BATCH_SIZE - 1)
#define GPU_OUTER_ITERS_PER_LAUNCH 40   // Batch consecutivi per thread per singolo kernel launch
#define BLOOM_SIZE_BITS 24              // 16 MB Bloom filter (ridotto per lasciare spazio ai buffer per-thread)

#define CUDA_CHECK(call) do { \
    cudaError_t err__ = (call); \
    if (err__ != cudaSuccess) { \
        fprintf(stderr, "[CUDA ERROR] %s:%d: %s\n", __FILE__, __LINE__, cudaGetErrorString(err__)); \
        exit(1); \
    } \
} while (0)

// ============================================================================
// ARITMETICA 256 BIT MOD p (secp256k1 field prime)
// limb[0] = 64 bit meno significativi ... limb[3] = 64 bit più significativi
// ============================================================================

struct u256 { uint64_t d[4]; };

__constant__ u256 FIELD_P = {{ 0xFFFFFFFEFFFFFC2FULL, 0xFFFFFFFFFFFFFFFFULL,
                                0xFFFFFFFFFFFFFFFFULL, 0xFFFFFFFFFFFFFFFFULL }};
// p - 2, usato per l'inversione modulare via piccolo teorema di Fermat (a^(p-2) mod p)
__constant__ u256 FIELD_P_MINUS_2 = {{ 0xFFFFFFFEFFFFFC2DULL, 0xFFFFFFFFFFFFFFFFULL,
                                        0xFFFFFFFFFFFFFFFFULL, 0xFFFFFFFFFFFFFFFFULL }};
// 2^256 mod p = 2^32 + 977, costante per la riduzione veloce (forma speciale del primo)
#define FIELD_K 0x1000003D1ULL

__host__ __device__ inline uint64_t load_be64(const uint8_t* p) {
    return ((uint64_t)p[0] << 56) | ((uint64_t)p[1] << 48) | ((uint64_t)p[2] << 40) | ((uint64_t)p[3] << 32) |
           ((uint64_t)p[4] << 24) | ((uint64_t)p[5] << 16) | ((uint64_t)p[6] << 8)  | ((uint64_t)p[7]);
}
__host__ __device__ inline void store_be64(uint8_t* p, uint64_t v) {
    p[0] = (uint8_t)(v >> 56); p[1] = (uint8_t)(v >> 48); p[2] = (uint8_t)(v >> 40); p[3] = (uint8_t)(v >> 32);
    p[4] = (uint8_t)(v >> 24); p[5] = (uint8_t)(v >> 16); p[6] = (uint8_t)(v >> 8);  p[7] = (uint8_t)(v);
}
__host__ __device__ inline u256 be32_to_u256(const uint8_t* b) {
    u256 r;
    r.d[3] = load_be64(b);
    r.d[2] = load_be64(b + 8);
    r.d[1] = load_be64(b + 16);
    r.d[0] = load_be64(b + 24);
    return r;
}
__host__ __device__ inline void u256_to_be32(const u256& a, uint8_t* b) {
    store_be64(b, a.d[3]);
    store_be64(b + 8, a.d[2]);
    store_be64(b + 16, a.d[1]);
    store_be64(b + 24, a.d[0]);
}

__device__ inline int u256_add_raw(const u256& a, const u256& b, u256& r) {
    unsigned __int128 carry = 0;
    for (int i = 0; i < 4; i++) {
        unsigned __int128 s = (unsigned __int128)a.d[i] + b.d[i] + carry;
        r.d[i] = (uint64_t)s;
        carry = s >> 64;
    }
    return (int)carry;
}
__device__ inline int u256_sub_raw(const u256& a, const u256& b, u256& r) {
    __int128 borrow = 0;
    for (int i = 0; i < 4; i++) {
        __int128 s = (__int128)a.d[i] - (__int128)b.d[i] - borrow;
        if (s < 0) { s += ((__int128)1 << 64); borrow = 1; } else borrow = 0;
        r.d[i] = (uint64_t)s;
    }
    return (int)borrow;
}
__device__ inline bool u256_ge(const u256& a, const u256& b) {
    for (int i = 3; i >= 0; i--) {
        if (a.d[i] != b.d[i]) return a.d[i] > b.d[i];
    }
    return true;
}

__device__ inline u256 addmod(const u256& a, const u256& b) {
    u256 r;
    int c = u256_add_raw(a, b, r);
    if (c || u256_ge(r, FIELD_P)) u256_sub_raw(r, FIELD_P, r);
    return r;
}
__device__ inline u256 submod(const u256& a, const u256& b) {
    u256 r;
    int borrow = u256_sub_raw(a, b, r);
    if (borrow) u256_add_raw(r, FIELD_P, r);
    return r;
}

// Moltiplicazione scolastica 256x256 -> 512 bit (8 limb), con propagazione
// del riporto corretta anche quando il riporto di una riga trabocca su più
// limb superiori (caso generale, non solo il limb immediatamente successivo).
__device__ inline void mul256_512(const u256& a, const u256& b, uint64_t r[8]) {
    for (int k = 0; k < 8; k++) r[k] = 0;
    for (int i = 0; i < 4; i++) {
        uint64_t carry = 0;
        for (int j = 0; j < 4; j++) {
            unsigned __int128 t = (unsigned __int128)a.d[i] * b.d[j] + r[i + j] + carry;
            r[i + j] = (uint64_t)t;
            carry = (uint64_t)(t >> 64);
        }
        int k = i + 4;
        while (carry) {
            unsigned __int128 t = (unsigned __int128)r[k] + carry;
            r[k] = (uint64_t)t;
            carry = (uint64_t)(t >> 64);
            k++;
        }
    }
}

// Riduzione mod p sfruttando 2^256 ≡ 2^32+977 (mod p): ogni "piega" sostituisce
// i bit oltre il 256-esimo con la loro immagine moltiplicata per FIELD_K,
// finché non resta più nulla da piegare; poi normalizza in [0, p).
__device__ inline u256 reduce512(const uint64_t Lin[8]) {
    uint64_t lo[4] = { Lin[0], Lin[1], Lin[2], Lin[3] };
    uint64_t hi[4] = { Lin[4], Lin[5], Lin[6], Lin[7] };

    while (hi[0] | hi[1] | hi[2] | hi[3]) {
        uint64_t t[5];
        uint64_t carry = 0;
        for (int i = 0; i < 4; i++) {
            unsigned __int128 p = (unsigned __int128)hi[i] * FIELD_K + carry;
            t[i] = (uint64_t)p;
            carry = (uint64_t)(p >> 64);
        }
        t[4] = carry;

        uint64_t newlo[4];
        uint64_t c = 0;
        for (int i = 0; i < 4; i++) {
            unsigned __int128 s = (unsigned __int128)lo[i] + t[i] + c;
            newlo[i] = (uint64_t)s;
            c = (uint64_t)(s >> 64);
        }
        unsigned __int128 ov = (unsigned __int128)t[4] + c;

        lo[0] = newlo[0]; lo[1] = newlo[1]; lo[2] = newlo[2]; lo[3] = newlo[3];
        hi[0] = (uint64_t)ov; hi[1] = (uint64_t)(ov >> 64); hi[2] = 0; hi[3] = 0;
    }

    u256 result = { { lo[0], lo[1], lo[2], lo[3] } };
    while (u256_ge(result, FIELD_P)) {
        u256 tmp;
        u256_sub_raw(result, FIELD_P, tmp);
        result = tmp;
    }
    return result;
}

__device__ inline u256 mulmod(const u256& a, const u256& b) {
    uint64_t t[8];
    mul256_512(a, b, t);
    return reduce512(t);
}

// Inversione modulare via piccolo teorema di Fermat: a^(p-2) mod p.
// Costosa (~256 quadrati + fino a 256 moltiplicazioni), usata una sola volta
// per batch (per invertire il prodotto totale nel trucco di Montgomery),
// quindi il costo è ammortizzato su GPU_EC_BATCH_SIZE chiavi.
__device__ inline u256 modinv_fermat(const u256& a) {
    u256 result = { { 1, 0, 0, 0 } };
    u256 base = a;
    for (int limb = 3; limb >= 0; limb--) {
        uint64_t e = FIELD_P_MINUS_2.d[limb];
        for (int bit = 63; bit >= 0; bit--) {
            result = mulmod(result, result);
            if ((e >> bit) & 1ULL) {
                result = mulmod(result, base);
            }
        }
    }
    return result;
}

// ============================================================================
// CURVA secp256k1 (a = 0): operazioni su punti Jacobiani
// ============================================================================

struct Jacobian { u256 X, Y, Z; };

// Raddoppio "dbl-2007-bl" specializzato per a=0
__device__ inline Jacobian jacobian_double(const Jacobian& P) {
    u256 A = mulmod(P.X, P.X);
    u256 B = mulmod(P.Y, P.Y);
    u256 C = mulmod(B, B);
    u256 xb = addmod(P.X, B);
    u256 D = submod(mulmod(xb, xb), addmod(A, C));
    D = addmod(D, D);
    u256 E = addmod(addmod(A, A), A); // 3*A
    u256 F = mulmod(E, E);
    Jacobian R;
    R.X = submod(F, addmod(D, D));
    u256 c8 = addmod(addmod(C, C), addmod(C, C));
    c8 = addmod(c8, c8); // 8*C
    R.Y = submod(mulmod(E, submod(D, R.X)), c8);
    u256 yz = addmod(P.Y, P.Z);
    R.Z = submod(mulmod(yz, yz), addmod(B, mulmod(P.Z, P.Z)));
    return R;
}

// Addizione mista Jacobiana + affine (Z2=1), formula generale "madd-2004-hmv"
// (Z1 qualsiasi). Usata per l'addizione di G durante la moltiplicazione
// scalare double-and-add. In caso di H=0 (collisione X, evento ~2^-256 per
// punti indipendenti) il risultato non è definito da queste formule: si
// marca Z3=0 per segnalare "punto da ignorare" (mai innescato in pratica).
__device__ inline Jacobian jacobian_add_mixed(const Jacobian& P, const u256& x2, const u256& y2) {
    u256 z1z1 = mulmod(P.Z, P.Z);
    u256 u2 = mulmod(x2, z1z1);
    u256 s2 = mulmod(mulmod(y2, P.Z), z1z1);
    u256 H = submod(u2, P.X);
    Jacobian R;
    if (H.d[0] == 0 && H.d[1] == 0 && H.d[2] == 0 && H.d[3] == 0) {
        R.Z = { { 0, 0, 0, 0 } };
        return R;
    }
    u256 HH = mulmod(H, H);
    u256 I = addmod(HH, HH); I = addmod(I, I); // 4*HH
    u256 J = mulmod(H, I);
    u256 r = addmod(submod(s2, P.Y), submod(s2, P.Y)); // 2*(S2-Y1)
    u256 V = mulmod(P.X, I);
    R.X = submod(submod(mulmod(r, r), J), addmod(V, V));
    u256 y1j = mulmod(P.Y, J);
    R.Y = submod(mulmod(r, submod(V, R.X)), addmod(y1j, y1j));
    u256 zh = addmod(P.Z, H);
    R.Z = submod(mulmod(zh, zh), addmod(z1z1, HH));
    return R;
}

// Addizione affine+affine con Z1=1 (usata SOLO per il batch: il punto base
// P0 è sempre stato appena normalizzato in affine). Stesse formule già
// validate nella versione CPU (vedi p2pk_bruteforce.cpp).
__device__ inline void jacobian_add_affine_z1(const u256& x1, const u256& y1,
                                               const u256& x2, const u256& y2,
                                               u256& X3, u256& Y3, u256& Z3) {
    u256 H = submod(x2, x1);
    if (H.d[0] == 0 && H.d[1] == 0 && H.d[2] == 0 && H.d[3] == 0) {
        Z3 = { { 0, 0, 0, 0 } };
        return;
    }
    u256 HH = mulmod(H, H);
    u256 HHH = mulmod(H, HH);
    u256 r = submod(y2, y1);
    X3 = submod(submod(mulmod(r, r), HHH), addmod(mulmod(x1, HH), mulmod(x1, HH)));
    Y3 = submod(mulmod(r, submod(mulmod(x1, HH), X3)), mulmod(y1, HHH));
    Z3 = H;
}

// Moltiplicazione scalare double-and-add (non ottimizzata: niente wNAF né
// tabelle precalcolate, a differenza della libsecp256k1 usata su CPU).
// Eseguita una sola volta ogni GPU_EC_BATCH_SIZE chiavi, quindi il suo costo
// (~256 raddoppi + ~128 addizioni in media) è ammortizzato sul batch.
__device__ inline Jacobian scalar_mult_basic(const u256& scalar_be_limbs, const u256& gx, const u256& gy) {
    Jacobian acc;
    acc.X = { { 0, 0, 0, 0 } };
    acc.Y = { { 0, 0, 0, 0 } };
    acc.Z = { { 0, 0, 0, 0 } }; // punto all'infinito
    bool started = false;

    for (int limb = 3; limb >= 0; limb--) {
        uint64_t e = scalar_be_limbs.d[limb];
        for (int bit = 63; bit >= 0; bit--) {
            if (started) acc = jacobian_double(acc);
            if ((e >> bit) & 1ULL) {
                if (!started) {
                    acc.X = gx; acc.Y = gy; acc.Z = { { 1, 0, 0, 0 } };
                    started = true;
                } else {
                    acc = jacobian_add_mixed(acc, gx, gy);
                }
            }
        }
    }
    return acc;
}

// ============================================================================
// GENERATORE PSEUDOCASUALE PER LA CHIAVE INIZIALE DI OGNI THREAD (xorshift128)
// ============================================================================

__device__ inline uint64_t xorshift64(uint64_t& s) {
    s ^= s << 13;
    s ^= s >> 7;
    s ^= s << 17;
    return s;
}

__device__ inline u256 random_start_privkey(uint64_t global_seed, uint32_t launch_id, uint32_t thread_id) {
    uint64_t s = global_seed ^ ((uint64_t)launch_id << 32) ^ ((uint64_t)thread_id * 0x9E3779B97F4A7C15ULL);
    if (s == 0) s = 0xDEADBEEFCAFEBABEULL;
    u256 r;
    r.d[0] = xorshift64(s);
    r.d[1] = xorshift64(s);
    r.d[2] = xorshift64(s);
    r.d[3] = xorshift64(s) & 0x7FFFFFFFFFFFFFFFULL; // resta ben sotto l'ordine della curva
    return r;
}

// Somma un intero piccolo (< 2^32) a uno scalare a 256 bit (limb little-endian)
__device__ inline void add_small_u256(u256& v, uint32_t n) {
    unsigned __int128 s = (unsigned __int128)v.d[0] + n;
    v.d[0] = (uint64_t)s;
    uint64_t carry = (uint64_t)(s >> 64);
    for (int i = 1; i < 4 && carry; i++) {
        s = (unsigned __int128)v.d[i] + carry;
        v.d[i] = (uint64_t)s;
        carry = (uint64_t)(s >> 64);
    }
}

// ============================================================================
// BLOOM FILTER + LOOKUP TARGET (device)
// ============================================================================

__device__ inline uint64_t bloom_hash1(const uint8_t* d) { const uint64_t* p = (const uint64_t*)d; return p[0] ^ (p[1] << 7); }
__device__ inline uint64_t bloom_hash2(const uint8_t* d) { const uint64_t* p = (const uint64_t*)d; return p[2] ^ (p[3] << 13); }
__device__ inline uint64_t bloom_hash3(const uint8_t* d) { const uint64_t* p = (const uint64_t*)d; return p[4] ^ (p[5] << 19); }

__device__ inline bool bloom_might_contain(const uint64_t* bits, uint64_t mask, const uint8_t* data) {
    uint64_t h1 = bloom_hash1(data) & mask;
    uint64_t h2 = bloom_hash2(data) & mask;
    uint64_t h3 = bloom_hash3(data) & mask;
    return (bits[h1 >> 6] & (1ULL << (h1 & 63))) &&
           (bits[h2 >> 6] & (1ULL << (h2 & 63))) &&
           (bits[h3 >> 6] & (1ULL << (h3 & 63)));
}

// Record ordinato per la ricerca binaria di conferma dopo un hit del Bloom filter
struct TargetRecord { uint8_t key[64]; uint32_t orig_idx; };

__device__ inline int bytes64_cmp(const uint8_t* a, const uint8_t* b) {
    for (int i = 0; i < 64; i++) {
        if (a[i] != b[i]) return (int)a[i] - (int)b[i];
    }
    return 0;
}

__device__ inline int find_target(const TargetRecord* sorted, int n, const uint8_t* key64) {
    int lo = 0, hi = n - 1;
    while (lo <= hi) {
        int mid = (lo + hi) / 2;
        int cmp = bytes64_cmp(sorted[mid].key, key64);
        if (cmp == 0) return (int)sorted[mid].orig_idx;
        if (cmp < 0) lo = mid + 1; else hi = mid - 1;
    }
    return -1;
}

__device__ inline int check_match(const uint64_t* bloom, uint64_t bloom_mask,
                                   const TargetRecord* sorted, int n_targets,
                                   const uint8_t* key64) {
    if (!bloom_might_contain(bloom, bloom_mask, key64)) return -1;
    return find_target(sorted, n_targets, key64);
}

// ============================================================================
// KERNEL DI RICERCA
// ============================================================================

__constant__ u256 D_GX;
__constant__ u256 D_GY;
__constant__ u256 D_PRECOMP_GX[GPU_EC_BATCH_MULT];
__constant__ u256 D_PRECOMP_GY[GPU_EC_BATCH_MULT];

__global__ void search_kernel(const uint64_t* bloom_bits, uint64_t bloom_mask,
                               const TargetRecord* sorted_targets, int n_targets,
                               uint64_t global_seed, uint32_t launch_id,
                               unsigned long long* d_total_attempts,
                               int* d_found_flag, uint8_t* d_found_privkey, int* d_found_target_idx) {
    uint32_t tid = blockIdx.x * blockDim.x + threadIdx.x;

    u256 privkey = random_start_privkey(global_seed, launch_id, tid);
    unsigned long long local_attempts = 0;

    u256 Xj[GPU_EC_BATCH_MULT], Yj[GPU_EC_BATCH_MULT], Zj[GPU_EC_BATCH_MULT];
    u256 invZ[GPU_EC_BATCH_MULT], prefix[GPU_EC_BATCH_MULT];
    bool valid[GPU_EC_BATCH_MULT];

    for (int outer = 0; outer < GPU_OUTER_ITERS_PER_LAUNCH; outer++) {
        if (*d_found_flag) break;

        Jacobian P0j = scalar_mult_basic(privkey, D_GX, D_GY);
        u256 zinv = modinv_fermat(P0j.Z);
        u256 zinv2 = mulmod(zinv, zinv);
        u256 zinv3 = mulmod(zinv2, zinv);
        u256 x0 = mulmod(P0j.X, zinv2);
        u256 y0 = mulmod(P0j.Y, zinv3);

        uint8_t key0[64];
        u256_to_be32(x0, key0);
        u256_to_be32(y0, key0 + 32);
        int m = check_match(bloom_bits, bloom_mask, sorted_targets, n_targets, key0);
        if (m >= 0 && atomicCAS(d_found_flag, 0, 1) == 0) {
            u256_to_be32(privkey, d_found_privkey);
            *d_found_target_idx = m;
        }

        int valid_count = 0;
        for (int i = 0; i < GPU_EC_BATCH_MULT; i++) {
            jacobian_add_affine_z1(x0, y0, D_PRECOMP_GX[i], D_PRECOMP_GY[i], Xj[i], Yj[i], Zj[i]);
            valid[i] = !(Zj[i].d[0] == 0 && Zj[i].d[1] == 0 && Zj[i].d[2] == 0 && Zj[i].d[3] == 0);
            if (valid[i]) valid_count++;
        }

        if (valid_count > 0) {
            int last = -1;
            for (int i = 0; i < GPU_EC_BATCH_MULT; i++) {
                if (!valid[i]) continue;
                prefix[i] = (last < 0) ? Zj[i] : mulmod(prefix[last], Zj[i]);
                last = i;
            }

            u256 inv_tmp = modinv_fermat(prefix[last]);

            for (int i = last; i >= 0; i--) {
                if (!valid[i]) continue;
                int prev = -1;
                for (int j = i - 1; j >= 0; j--) { if (valid[j]) { prev = j; break; } }
                invZ[i] = (prev >= 0) ? mulmod(inv_tmp, prefix[prev]) : inv_tmp;
                inv_tmp = mulmod(inv_tmp, Zj[i]);
            }

            for (int i = 0; i < GPU_EC_BATCH_MULT; i++) {
                if (!valid[i]) continue;
                u256 iv2 = mulmod(invZ[i], invZ[i]);
                u256 iv3 = mulmod(iv2, invZ[i]);
                u256 ax = mulmod(Xj[i], iv2);
                u256 ay = mulmod(Yj[i], iv3);

                uint8_t key[64];
                u256_to_be32(ax, key);
                u256_to_be32(ay, key + 32);

                int mi = check_match(bloom_bits, bloom_mask, sorted_targets, n_targets, key);
                if (mi >= 0 && atomicCAS(d_found_flag, 0, 1) == 0) {
                    u256 found = privkey;
                    add_small_u256(found, (uint32_t)(i + 1));
                    u256_to_be32(found, d_found_privkey);
                    *d_found_target_idx = mi;
                }
            }
        }

        local_attempts += GPU_EC_BATCH_SIZE;
        add_small_u256(privkey, GPU_EC_BATCH_SIZE);
    }

    atomicAdd(d_total_attempts, local_attempts);
}

// ============================================================================
// HOST: caricamento target, precompute, main loop
// ============================================================================

struct TargetKey { uint8_t pubkey[65]; char hex[131]; };
static std::vector<TargetKey> g_target_keys;

static volatile int g_keep_running = 1;
static void sigint_handler(int) {
    g_keep_running = 0;
    printf("\n\n[!] Interruzione rilevata, chiusura in corso...\n");
}

static int hex_to_bytes(const char* hex, uint8_t* bytes, size_t len) {
    if (strlen(hex) != len * 2) return 0;
    for (size_t i = 0; i < len; i++) sscanf(hex + i * 2, "%2hhx", &bytes[i]);
    return 1;
}
static void bytes_to_hex(const uint8_t* bytes, size_t len, char* hex) {
    for (size_t i = 0; i < len; i++) sprintf(hex + i * 2, "%02x", bytes[i]);
    hex[len * 2] = '\0';
}

static void format_number(uint64_t num, char* buffer) {
    if (num >= 1000000000000ULL) sprintf(buffer, "%.2fT", num / 1000000000000.0);
    else if (num >= 1000000000ULL) sprintf(buffer, "%.2fG", num / 1000000000.0);
    else if (num >= 1000000ULL) sprintf(buffer, "%.2fM", num / 1000000.0);
    else if (num >= 1000ULL) sprintf(buffer, "%.2fK", num / 1000.0);
    else sprintf(buffer, "%lu", (unsigned long)num);
}

int main(int argc, char** argv) {
    printf("========================================\n");
    printf("  Bitcoin P2PK Bruteforce - GPU (CUDA)\n");
    printf("  SOLO PER SCOPI EDUCATIVI\n");
    printf("========================================\n\n");

    const char* target_file = "target_keys.txt";
    if (argc > 1) target_file = argv[1];

    int device_count = 0;
    CUDA_CHECK(cudaGetDeviceCount(&device_count));
    if (device_count == 0) {
        fprintf(stderr, "[ERROR] Nessuna GPU CUDA trovata\n");
        return 1;
    }
    cudaDeviceProp prop;
    CUDA_CHECK(cudaGetDeviceProperties(&prop, 0));
    printf("[+] GPU: %s (SM %d.%d, %d multiprocessori, %.1f GB)\n",
           prop.name, prop.major, prop.minor, prop.multiProcessorCount,
           prop.totalGlobalMem / 1e9);

    secp256k1_context* ctx = secp256k1_context_create(SECP256K1_CONTEXT_SIGN | SECP256K1_CONTEXT_VERIFY);

    // Generatore secp256k1 (coordinate standard)
    uint8_t gx_bytes[32], gy_bytes[32];
    {
        uint8_t one[32] = {0}; one[31] = 1;
        secp256k1_pubkey g_pub;
        if (!secp256k1_ec_pubkey_create(ctx, &g_pub, one)) { fprintf(stderr, "[ERROR] pubkey_create(G)\n"); return 1; }
        unsigned char buf[65]; size_t outlen = 65;
        secp256k1_ec_pubkey_serialize(ctx, buf, &outlen, &g_pub, SECP256K1_EC_UNCOMPRESSED);
        memcpy(gx_bytes, buf + 1, 32);
        memcpy(gy_bytes, buf + 33, 32);
    }
    u256 h_gx = be32_to_u256(gx_bytes);
    u256 h_gy = be32_to_u256(gy_bytes);
    CUDA_CHECK(cudaMemcpyToSymbol(D_GX, &h_gx, sizeof(u256)));
    CUDA_CHECK(cudaMemcpyToSymbol(D_GY, &h_gy, sizeof(u256)));

    // Multipli precalcolati 1G..GPU_EC_BATCH_MULT*G
    printf("[+] Precalcolo multipli del generatore (1G..%dG)...\n", GPU_EC_BATCH_MULT);
    {
        std::vector<u256> h_pgx(GPU_EC_BATCH_MULT), h_pgy(GPU_EC_BATCH_MULT);
        for (int i = 0; i < GPU_EC_BATCH_MULT; i++) {
            uint8_t pk[32] = {0};
            uint32_t v = (uint32_t)(i + 1);
            pk[31] = v & 0xFF; pk[30] = (v >> 8) & 0xFF; pk[29] = (v >> 16) & 0xFF;
            secp256k1_pubkey pub;
            if (!secp256k1_ec_pubkey_create(ctx, &pub, pk)) { fprintf(stderr, "[ERROR] precompute %dG\n", i + 1); return 1; }
            unsigned char buf[65]; size_t outlen = 65;
            secp256k1_ec_pubkey_serialize(ctx, buf, &outlen, &pub, SECP256K1_EC_UNCOMPRESSED);
            h_pgx[i] = be32_to_u256(buf + 1);
            h_pgy[i] = be32_to_u256(buf + 33);
        }
        CUDA_CHECK(cudaMemcpyToSymbol(D_PRECOMP_GX, h_pgx.data(), sizeof(u256) * GPU_EC_BATCH_MULT));
        CUDA_CHECK(cudaMemcpyToSymbol(D_PRECOMP_GY, h_pgy.data(), sizeof(u256) * GPU_EC_BATCH_MULT));
    }

    // Caricamento target keys (stesso formato/euristica della versione CPU)
    printf("[+] Caricamento chiavi target da %s...\n", target_file);
    std::vector<TargetRecord> h_sorted;
    {
        std::ifstream file(target_file);
        if (!file.is_open()) { fprintf(stderr, "[ERROR] Impossibile aprire %s\n", target_file); return 1; }
        std::string line;
        std::getline(file, line); // header
        while (std::getline(file, line)) {
            if (line.empty()) continue;
            std::string hexs = line;
            hexs.erase(remove_if(hexs.begin(), hexs.end(), [](unsigned char c){ return isspace(c); }), hexs.end());
            if (hexs.length() != 130 && hexs.length() != 128) continue;
            if (hexs.length() == 128) hexs = "04" + hexs;

            TargetKey key;
            if (!hex_to_bytes(hexs.c_str(), key.pubkey, 65)) continue;
            strcpy(key.hex, hexs.c_str());

            secp256k1_pubkey pub;
            if (!secp256k1_ec_pubkey_parse(ctx, &pub, key.pubkey, 65)) continue;

            TargetRecord rec;
            memcpy(rec.key, key.pubkey + 1, 64);
            rec.orig_idx = (uint32_t)g_target_keys.size();
            g_target_keys.push_back(key);
            h_sorted.push_back(rec);
        }
    }
    if (h_sorted.empty()) { fprintf(stderr, "[ERROR] Nessuna chiave target caricata\n"); return 1; }
    std::sort(h_sorted.begin(), h_sorted.end(), [](const TargetRecord& a, const TargetRecord& b) {
        return memcmp(a.key, b.key, 64) < 0;
    });
    printf("[+] Caricate %zu chiavi pubbliche target\n", g_target_keys.size());

    TargetRecord* d_sorted_targets;
    CUDA_CHECK(cudaMalloc(&d_sorted_targets, sizeof(TargetRecord) * h_sorted.size()));
    CUDA_CHECK(cudaMemcpy(d_sorted_targets, h_sorted.data(), sizeof(TargetRecord) * h_sorted.size(), cudaMemcpyHostToDevice));

    // Bloom filter
    uint64_t bloom_words = (1ULL << BLOOM_SIZE_BITS) / 64;
    uint64_t bloom_mask = (1ULL << BLOOM_SIZE_BITS) - 1;
    std::vector<uint64_t> h_bloom(bloom_words, 0);
    auto bloom_add = [&](const uint8_t* d) {
        const uint64_t* p = (const uint64_t*)d;
        uint64_t h1 = (p[0] ^ (p[1] << 7)) & bloom_mask;
        uint64_t h2 = (p[2] ^ (p[3] << 13)) & bloom_mask;
        uint64_t h3 = (p[4] ^ (p[5] << 19)) & bloom_mask;
        h_bloom[h1 >> 6] |= (1ULL << (h1 & 63));
        h_bloom[h2 >> 6] |= (1ULL << (h2 & 63));
        h_bloom[h3 >> 6] |= (1ULL << (h3 & 63));
    };
    for (auto& r : h_sorted) bloom_add(r.key);
    uint64_t* d_bloom;
    CUDA_CHECK(cudaMalloc(&d_bloom, sizeof(uint64_t) * bloom_words));
    CUDA_CHECK(cudaMemcpy(d_bloom, h_bloom.data(), sizeof(uint64_t) * bloom_words, cudaMemcpyHostToDevice));
    printf("[+] Bloom filter: %llu MB\n", (unsigned long long)(bloom_words * 8 / 1024 / 1024));

    // Buffer risultati
    unsigned long long* d_total_attempts;
    int* d_found_flag;
    uint8_t* d_found_privkey;
    int* d_found_target_idx;
    CUDA_CHECK(cudaMalloc(&d_total_attempts, sizeof(unsigned long long)));
    CUDA_CHECK(cudaMalloc(&d_found_flag, sizeof(int)));
    CUDA_CHECK(cudaMalloc(&d_found_privkey, 32));
    CUDA_CHECK(cudaMalloc(&d_found_target_idx, sizeof(int)));
    CUDA_CHECK(cudaMemset(d_total_attempts, 0, sizeof(unsigned long long)));
    CUDA_CHECK(cudaMemset(d_found_flag, 0, sizeof(int)));

    int threads_per_block = 256;
    int blocks = prop.multiProcessorCount * 8; // occupazione euristica, non ottimizzata a fondo
    uint64_t total_threads = (uint64_t)threads_per_block * blocks;
    printf("[+] Griglia CUDA: %d blocchi x %d thread = %llu thread totali\n",
           blocks, threads_per_block, (unsigned long long)total_threads);
    printf("[+] Batch per thread: %d chiavi | Batch per kernel launch: %d\n\n",
           GPU_EC_BATCH_SIZE, GPU_OUTER_ITERS_PER_LAUNCH);

    signal(SIGINT, sigint_handler);
    signal(SIGTERM, sigint_handler);

    time_t start_time = time(NULL);
    uint64_t last_total = 0;
    struct timespec last_ts; clock_gettime(CLOCK_MONOTONIC, &last_ts);
    uint32_t launch_id = 0;
    uint64_t global_seed = (uint64_t)start_time ^ 0xA5A5A5A5DEADBEEFULL;

    int found = 0;
    while (g_keep_running && !found) {
        search_kernel<<<blocks, threads_per_block>>>(
            d_bloom, bloom_mask, d_sorted_targets, (int)h_sorted.size(),
            global_seed, launch_id++, d_total_attempts, d_found_flag, d_found_privkey, d_found_target_idx);
        CUDA_CHECK(cudaGetLastError());
        CUDA_CHECK(cudaDeviceSynchronize());

        unsigned long long total;
        CUDA_CHECK(cudaMemcpy(&total, d_total_attempts, sizeof(unsigned long long), cudaMemcpyDeviceToHost));
        CUDA_CHECK(cudaMemcpy(&found, d_found_flag, sizeof(int), cudaMemcpyDeviceToHost));

        struct timespec now_ts; clock_gettime(CLOCK_MONOTONIC, &now_ts);
        double window_sec = (now_ts.tv_sec - last_ts.tv_sec) + (now_ts.tv_nsec - last_ts.tv_nsec) / 1e9;
        if (window_sec < 0.001) window_sec = 0.001;
        double rate = (total - last_total) / window_sec;
        double elapsed_total = difftime(time(NULL), start_time);
        if (elapsed_total < 1) elapsed_total = 1;

        char total_str[32], rate_str[32];
        format_number(total, total_str);
        format_number((uint64_t)rate, rate_str);
        printf("[INFO] Tentativi totali: %s | Velocità: %s keys/sec (GPU) | Tempo: %.0fs\n",
               total_str, rate_str, elapsed_total);

        last_total = total;
        last_ts = now_ts;
    }

    if (found) {
        uint8_t priv[32];
        int target_idx;
        CUDA_CHECK(cudaMemcpy(priv, d_found_privkey, 32, cudaMemcpyDeviceToHost));
        CUDA_CHECK(cudaMemcpy(&target_idx, d_found_target_idx, sizeof(int), cudaMemcpyDeviceToHost));

        char priv_hex[65];
        bytes_to_hex(priv, 32, priv_hex);

        printf("\n\n========================================\n");
        printf("🎯 CHIAVE TROVATA! 🎯\n");
        printf("========================================\n");
        printf("Private Key: %s\n", priv_hex);
        printf("Public Key:  %s\n", g_target_keys[target_idx].hex);
        printf("========================================\n\n");

        FILE* f = fopen("found_keys.txt", "a");
        if (f) {
            time_t now = time(NULL);
            fprintf(f, "\n=== FOUND (GPU) at %s", ctime(&now));
            fprintf(f, "Private Key: %s\n", priv_hex);
            fprintf(f, "Public Key:  %s\n", g_target_keys[target_idx].hex);
            fprintf(f, "========================================\n");
            fclose(f);
        }
    }

    cudaFree(d_bloom);
    cudaFree(d_sorted_targets);
    cudaFree(d_total_attempts);
    cudaFree(d_found_flag);
    cudaFree(d_found_privkey);
    cudaFree(d_found_target_idx);
    secp256k1_context_destroy(ctx);

    printf("[+] Programma terminato\n");
    return 0;
}