Aggiungi versione GPU (CUDA) del bruteforce

Porta la stessa strategia algoritmica della versione CPU (batch EC in
coordinate Jacobiane con Z1=1 + un'unica inversione di campo per batch
via trucco di Montgomery) su GPU NVIDIA, dove però GMP non è
disponibile: l'aritmetica a 256 bit mod il primo di secp256k1 è scritta
a mano (limb a 64 bit + __int128, riduzione veloce sfruttando
2^256 ≡ 2^32+977 mod p), e ogni thread CUDA agisce come una lane di
ricerca indipendente con la propria chiave privata iniziale casuale —
esattamente come un thread CPU, ma migliaia in parallelo invece di una
manciata.

La moltiplicazione scalare su device è un semplice double-and-add
(nessun wNAF/tabelle precalcolate come in libsecp256k1 lato host), il
principale margine di ottimizzazione futura. Il matching usa un Bloom
filter + ricerca binaria su un array di target ordinato, entrambi
caricati una volta sola in memoria device.

Aggiunge i target `make gpu` / `make gpu-info` al Makefile (NVCC_ARCH
configurabile, default sm_61/Pascal) e documenta in README che su WSL2
il driver NVIDIA arriva dall'host Windows: dentro WSL serve installare
solo il CUDA Toolkit, mai un driver.

Verificato: aritmetica di campo, moltiplicazione scalare e batch add
confrontati bit-per-bit con secp256k1_ec_pubkey_create (nessuna
discrepanza su ~30 valori di test); run end-to-end con privkey nota
trova la chiave corretta; nessun falso positivo su chiavi pubbliche
valide casuali. Su una Quadro P4000 (Pascal): ~15M keys/sec, contro
~3.5-4M keys/sec della CPU con 11 thread nello stesso ambiente.

.gitignore

@@ -29,6 +29,7 @@ build/
 # Eseguibili compilati
 bruteforce/p2pk_bruteforce
 bruteforce/p2pk_bruteforce_debug
+bruteforce/p2pk_bruteforce_gpu
 
 # File oggetto e compilazione
 *.o

CLAUDE.md

@@ -0,0 +1,84 @@
+# CLAUDE.md
+
+This file provides guidance to Claude Code (claude.ai/code) when working with code in this repository.
+
+## Project purpose
+
+Educational/research suite for studying Bitcoin P2PK (Pay-to-Public-Key) transactions and demonstrating why ECDSA secp256k1 bruteforce is computationally infeasible (keyspace 2^256). All docs and CLI output are in Italian. Two independent components share data via files, not code:
+
+1. `databases/` — Python scanner that walks the Bitcoin blockchain via the mempool.space API, finds P2PK outputs, stores them in SQLite, and checks UTXO spent/unspent status.
+2. `bruteforce/` — CPU (C++/pthreads) and GPU (CUDA) programs that load target public keys and search the private-key space for matches, as a performance demonstration (not a realistic attack — success probability is ~2^-256).
+
+Data flows one way: scanner → SQLite DB (`databases/bitcoin_p2pk_study.db`) → `extract_p2pk_utxo.py` filters unspent P2PK → `target_keys.txt` → C++ bruteforce consumes it.
+
+## Commands
+
+### Python scanner (`databases/`)
+```bash
+python3 -m venv .venv && source .venv/bin/activate
+pip install -r requirements.txt
+
+cd databases
+python3 scan_blockchain.py   # interactive: prompts for start/end block + delay
+python3 view_db.py           # generates p2pk_report.html
+python3 view_db.py --stats   # prints stats to terminal
+```
+
+### C++ bruteforce (`bruteforce/`)
+```bash
+cd bruteforce
+make install-deps      # apt packages: build-essential, libsecp256k1-dev, libgmp-dev, autoconf, libtool, pkg-config
+make                   # builds secp256k1 locally from source if bruteforce/secp256k1/ doesn't exist (~5 min first time), then compiles p2pk_bruteforce
+make clean             # remove binaries/object files
+make clean-all         # also removes the locally-built secp256k1 tree
+make debug             # -g -O0 build: p2pk_bruteforce_debug
+make bench             # 10-second timed run
+make pgo               # profile-guided optimization build (3-step: generate → run 30s → use)
+make valgrind          # leak check on the debug build
+make help              # list all targets
+
+python3 extract_p2pk_utxo.py [db_path] [output.txt]   # default: ../databases/bitcoin_p2pk_study.db -> target_keys.txt
+python3 extract_p2pk_utxo.py --stats                  # DB stats only, no extraction
+
+./p2pk_bruteforce [target_keys.txt]   # runs until Ctrl+C; logs to progress.csv, matches to found_keys.txt
+
+# GPU version (CUDA), much faster if an NVIDIA GPU + driver + CUDA Toolkit are available
+make gpu-info          # checks nvidia-smi and nvcc are present
+make gpu                # NVCC_ARCH defaults to sm_61 (Pascal) — override for other GPUs, e.g. make gpu NVCC_ARCH=sm_86
+./p2pk_bruteforce_gpu [target_keys.txt]
+```
+
+There is no test suite in this repo.
+
+## Architecture notes
+
+### `bruteforce/p2pk_bruteforce.cpp` (CPU, single file, everything lives here)
+- One pthread per (core count − 1); each thread gets a disjoint slice of the 256-bit keyspace via `partition_keyspace` (only the top 64 bits are partitioned — the search relies on random start offsets within each thread's slice, not a full 256-bit range split).
+- Target pubkeys are loaded from a `.txt` file (uncompressed hex, `04` prefix, one per line, header line skipped) into both an `unordered_map<std::array<uint8_t,64>,...>` keyed on raw X||Y bytes (exact match) and a 64MB Bloom filter (fast negative rejection, checked first via `check_match_fast_raw`).
+- Performance trick: one EC scalar multiplication (via libsecp256k1) produces a base point P0, then the next `EC_BATCH_SIZE` (256) keys are derived via Jacobian-coordinate EC point addition (`ec_add_affine_affine`, formulas specialized for Z1=1) against precomputed multiples of G, converted back to affine with a single shared modular inversion per batch (Montgomery's batch-inversion trick, via GMP) instead of one inversion per key.
+- `increment_privkey`/`add_to_privkey` treat the 32-byte scalar as 4 native 64-bit words (little-endian on x86), not as one big-endian integer — this still enumerates the keyspace without collisions, just in a "scrambled" order; don't reuse them to reconstruct an exact scalar from EC math (use `add_small_be256` for that, as `save_found_key` does).
+- Build system auto-detects a local `bruteforce/secp256k1/` (built by `build_secp256k1.sh` targeting this specific CPU with `-march=native`) and links against it via rpath instead of the system lib; falls back to system `libsecp256k1`/`libgmp` if absent.
+- Matches are written to `found_keys.txt`; throughput stats (instantaneous rate, summed across threads) print to stdout and `progress.csv` every `PROGRESS_INTERVAL_SEC` (2s).
+
+### `bruteforce/p2pk_bruteforce_gpu.cu` (GPU, CUDA, single file)
+- Same algorithmic strategy as the CPU version (Jacobian batch add + Montgomery batch inversion), but GMP doesn't run on-device, so 256-bit field arithmetic mod the secp256k1 prime is hand-written (`u256` = 4×uint64 limbs, schoolbook multiply via `unsigned __int128`, fast reduction using `2^256 ≡ 2^32+977 (mod p)`). Verified bit-for-bit against `secp256k1_ec_pubkey_create` before trusting it (see conversation/test methodology — no separate test file is checked in).
+- Each CUDA thread is an independent search lane (own random starting privkey, own local batch state), same as a CPU thread — just thousands of them instead of ~11. Per-thread batch size is `GPU_EC_BATCH_SIZE` (128), smaller than the CPU's 256, to keep per-thread local-memory footprint (Jacobian batch arrays) bounded across tens of thousands of concurrent threads.
+- Scalar multiplication on-device is plain double-and-add (no wNAF/windowing/precomputed tables like libsecp256k1 has on CPU) — the main further-optimization opportunity if more speed is needed.
+- Host side still uses libsecp256k1 (CPU) only for one-time setup: precomputing G-multiples and validating/loading target keys into a host-sorted array + Bloom filter, both uploaded once to device memory. Device-side matching = Bloom filter check + binary search over the sorted target array.
+- `NVCC_ARCH` in the Makefile defaults to `sm_61` (Pascal) — must match the actual GPU (`make gpu-info` shows compute capability via `nvidia-smi`, then pick the matching `sm_XX`).
+- On WSL2, the NVIDIA driver comes from the Windows host (visible as `nvidia-smi` working out of the box) — only the CUDA Toolkit needs installing inside WSL, never a driver.
+
+### `databases/scan_blockchain.py`
+- `P2PKBlockchainScanner` class wraps all mempool.space API access (`get_block_hash`, `get_block_transactions` with pagination, `check_utxo_status`) and SQLite persistence.
+- P2PK detection is deliberately redundant — a script is classified as P2PK if ANY of 4 independent checks match: explicit `scriptpubkey_type`, script byte length (67 or 35 bytes), ASM pattern (`<pubkey> OP_CHECKSIG`), or raw hex pattern (`41<pubkey>ac` / `21<pubkey>ac`). This exists because the API's `scriptpubkey_type` field is not reliable for very old (pre-2012) P2PK outputs.
+- Scanning is resumable: `scan_progress` table (single row, `id=1`) tracks `last_scanned_block`; reruns default to `last_scanned_block + 1`. `UNIQUE(txid, output_index)` on `p2pk_addresses` prevents duplicate inserts, so overlapping scan ranges are safe.
+- This script and its SQLite DB/CSV outputs are intended to be committed and shared across contributors scanning different block ranges (see `.gitignore` — DB/CSV/HTML are NOT excluded).
+
+### `bruteforce/extract_p2pk_utxo.py`
+- Reads only `is_unspent = 1` rows from the scanner's DB, strips the `41.../21...ac` script wrapper to get the raw pubkey, and re-adds the `04` prefix.
+- Compressed pubkeys (33-byte, script length 70/hex prefix `21`) are explicitly skipped — the C++ bruteforce only generates and matches uncompressed public keys.
+
+## Key coupling to be aware of
+
+- The bruteforce binary's target file format (uncompressed hex pubkeys, `04` prefix, header line) is produced exclusively by `extract_p2pk_utxo.py` — if editing one side's format, update the other.
+- `Makefile` CFLAGS use `-march=native`/`-mtune=native` and `-ffast-math`; binaries are not portable across different CPUs and should be rebuilt (`make clean && make`) after moving to different hardware.

bruteforce/Makefile

@@ -27,6 +27,14 @@ SOURCE = p2pk_bruteforce.cpp
 INCLUDE_PATH = -I/usr/local/include -I/usr/include
 LIB_PATH = -L/usr/local/lib -L/usr/lib
 
+# Versione GPU (CUDA)
+NVCC = nvcc
+# Architettura GPU target: sm_61 = Pascal (es. Quadro/GTX 10xx). Sovrascrivi con
+# `make gpu NVCC_ARCH=sm_75` (Turing), sm_86 (Ampere), sm_89 (Ada) etc. secondo la tua GPU.
+NVCC_ARCH = sm_61
+GPU_TARGET = p2pk_bruteforce_gpu
+GPU_SOURCE = p2pk_bruteforce_gpu.cu
+
 # ============================================================================
 # TARGET PRINCIPALI
 # ============================================================================
@@ -140,7 +148,7 @@ valgrind: debug
 # Pulizia
 clean:
 	@echo "[+] Pulizia file compilati..."
-	rm -f $(TARGET) $(TARGET)_debug $(TARGET)_pgo
+	rm -f $(TARGET) $(TARGET)_debug $(TARGET)_pgo $(GPU_TARGET)
 	rm -f *.o *.gcda *.gcno *.s
 	rm -f progress.csv
 	@echo "[+] Pulizia completata!"
@@ -170,6 +178,35 @@ build-secp256k1:
 	@echo "[+] Compilazione libsecp256k1 ottimizzata..."
 	@./build_secp256k1.sh
 
+# ============================================================================
+# VERSIONE GPU (CUDA)
+# ============================================================================
+
+# Richiede: NVIDIA CUDA Toolkit (nvcc) + driver NVIDIA funzionante (verifica
+# con `nvidia-smi`). Su WSL2 il driver è quello del sistema Windows host: non
+# va installato dentro WSL, basta il CUDA Toolkit (`nvcc --version` per controllare).
+gpu: $(GPU_SOURCE)
+	@echo "========================================="
+	@echo "  Bitcoin P2PK Bruteforce GPU - Compilazione"
+	@echo "========================================="
+	@which $(NVCC) > /dev/null || (echo "[ERROR] nvcc non trovato. Installa il CUDA Toolkit."; exit 1)
+	$(NVCC) -O3 -arch=$(NVCC_ARCH) -std=c++17 -diag-suppress 1650 \
+		$(INCLUDE_PATH) $(LIB_PATH) \
+		-o $(GPU_TARGET) $(GPU_SOURCE) -lsecp256k1
+	@echo "[+] Compilazione completata!"
+	@echo "[+] Eseguibile: ./$(GPU_TARGET)"
+	@echo ""
+	@echo "OTTIMIZZAZIONI ATTIVE:"
+	@echo "  ✓ Batch EC point addition in coordinate Jacobiane (stesso algoritmo della CPU)"
+	@echo "  ✓ Batch modular inversion (Montgomery) per thread"
+	@echo "  ✓ Migliaia di thread paralleli invece di N core CPU"
+	@echo "========================================="
+
+gpu-info:
+	@nvidia-smi --query-gpu=name,driver_version,memory.total,compute_cap --format=csv 2>&1 || \
+		echo "[!] nvidia-smi non disponibile: driver NVIDIA non rilevato"
+	@which $(NVCC) > /dev/null && $(NVCC) --version || echo "[!] nvcc non trovato: CUDA Toolkit non installato"
+
 # ============================================================================
 # HELP
 # ============================================================================
@@ -191,15 +228,22 @@ help:
 	@echo "  make clean-all          - Pulizia completa"
 	@echo "  make install-deps       - Installa dipendenze"
 	@echo "  make build-secp256k1    - Compila libsecp256k1 locale"
+	@echo "  make gpu                - Compila la versione GPU (richiede CUDA Toolkit)"
+	@echo "  make gpu-info           - Controlla driver NVIDIA e CUDA Toolkit disponibili"
 	@echo ""
 	@echo "Uso consigliato:"
 	@echo "  1. make                 # Compila"
 	@echo "  2. ./$(TARGET)          # Esegui bruteforce"
 	@echo ""
-	@echo "Per massime performance:"
+	@echo "Per massime performance su CPU:"
 	@echo "  make pgo                # Compila con PGO (+10-20% speed)"
 	@echo ""
+	@echo "Per usare la GPU (molto più veloce, se disponibile):"
+	@echo "  make gpu-info           # Verifica driver/CUDA Toolkit"
+	@echo "  make gpu                # Compila"
+	@echo "  ./$(GPU_TARGET)         # Esegui"
+	@echo ""
 	@echo "==================================================="
 
 .PHONY: all build pgo pgo-generate pgo-run pgo-use debug asm bench \
-        valgrind clean clean-all install-deps build-secp256k1 help
+        valgrind clean clean-all install-deps build-secp256k1 gpu gpu-info help

bruteforce/README.md

@@ -54,6 +54,32 @@ make
 
 ---
 
+## Versione GPU (CUDA)
+
+Se hai una GPU NVIDIA con driver e CUDA Toolkit disponibili, la versione GPU
+usa la stessa strategia algoritmica di quella CPU (batch EC in coordinate
+Jacobiane + una sola inversione di campo per batch) ma su migliaia di thread
+paralleli invece che su una manciata di core.
+
+```bash
+cd bruteforce
+make gpu-info    # verifica driver NVIDIA (nvidia-smi) e CUDA Toolkit (nvcc)
+make gpu         # compila (default: sm_61 / Pascal — sovrascrivi con NVCC_ARCH per altre GPU)
+./p2pk_bruteforce_gpu
+```
+
+Su una Quadro P4000 (Pascal, 2016) si osservano ~15M keys/sec, contro i
+~300K-2M keys/sec della CPU — ma la moltiplicazione scalare su GPU non è
+ottimizzata quanto quella di libsecp256k1 (niente wNAF/tabelle precalcolate),
+quindi il margine di ulteriore miglioramento è ancora ampio su GPU più
+recenti.
+
+> **Nota (WSL2)**: il driver NVIDIA arriva dal sistema Windows host (basta
+> che `nvidia-smi` funzioni già) — dentro WSL serve installare solo il CUDA
+> Toolkit, mai un driver.
+
+---
+
 ## Utilizzo
 
 ### 1. Prepara il File Target

bruteforce/p2pk_bruteforce.cpp

@@ -3,11 +3,14 @@
  * Versione CPU ottimizzata per massime prestazioni
  *
  * OTTIMIZZAZIONI IMPLEMENTATE:
- * - Batch EC point addition (genera N chiavi con 1 moltiplicazione + N addizioni)
- * - Zero-copy: niente serializzazione fino al match
- * - Hash diretto su secp256k1_pubkey raw data
+ * - Batch EC point addition in coordinate Jacobiane (1 moltiplicazione scalare
+ *   + N addizioni EC affine+affine, tutte con Z1=1) invece di N chiamate a
+ *   secp256k1_ec_pubkey_combine
+ * - Batch modular inversion (trucco di Montgomery): 1 sola inversione di campo
+ *   per batch di N chiavi invece di N inversioni separate
+ * - Zero-copy: matching diretto sui byte grezzi X||Y, niente oggetti
+ *   secp256k1_pubkey nel loop caldo
  * - SIMD-friendly Bloom filter
- * - Precomputed lookup tables
  * - Cache-aligned memory
  * - CPU prefetching hints
  *
@@ -22,9 +25,11 @@
 #include <signal.h>
 #include <unistd.h>
 #include <secp256k1.h>
+#include <gmp.h>
 #include <pthread.h>
 #include <sys/time.h>
 #include <vector>
+#include <array>
 #include <string>
 #include <unordered_map>
 #include <fstream>
@@ -41,12 +46,14 @@
 // CONFIGURAZIONE OTTIMIZZATA
 // ============================================================================
 
-#define EC_BATCH_SIZE 256      // Genera 256 chiavi consecutive con EC addition (+25% speed)
-#define SYNC_BATCH 100000      // Sincronizza contatori ogni 100K chiavi
+#define EC_BATCH_SIZE 256      // Genera 256 chiavi consecutive con EC addition
+#define EC_BATCH_MULT (EC_BATCH_SIZE - 1) // Multipli di G precalcolati necessari (2G..256G)
+#define SYNC_BATCH 50000       // Sincronizza contatori ogni 50K chiavi (aggiornamento terminale fluido)
 #define MAX_THREADS 256
 #define BLOOM_SIZE_BITS 26     // 64MB Bloom filter
 #define USE_BLOOM_FILTER 1
 #define USE_EC_BATCH 1         // Abilita batch EC point addition
+#define PROGRESS_INTERVAL_SEC 2 // Intervallo di aggiornamento statistiche a terminale
 
 // ============================================================================
 // STRUTTURE DATI OTTIMIZZATE
@@ -58,18 +65,17 @@ struct TargetKey {
     char hex[131];
 };
 
-// Hash ottimizzato per raw secp256k1_pubkey data (64 bytes)
-struct PubkeyRawHash {
-    size_t operator()(const secp256k1_pubkey& key) const {
-        const uint64_t* p = reinterpret_cast<const uint64_t*>(key.data);
-        // XOR rapido dei primi 64 bit
+// Hash ottimizzato per chiave grezza X||Y (64 byte, formato non compresso senza il prefisso 04)
+struct RawKeyHash {
+    size_t operator()(const std::array<uint8_t, 64>& key) const {
+        const uint64_t* p = reinterpret_cast<const uint64_t*>(key.data());
         return p[0] ^ p[1] ^ p[2];
     }
 };
 
-struct PubkeyRawEqual {
-    bool operator()(const secp256k1_pubkey& a, const secp256k1_pubkey& b) const {
-        return memcmp(a.data, b.data, 64) == 0;
+struct RawKeyEqual {
+    bool operator()(const std::array<uint8_t, 64>& a, const std::array<uint8_t, 64>& b) const {
+        return memcmp(a.data(), b.data(), 64) == 0;
     }
 };
 
@@ -82,7 +88,7 @@ private:
     size_t size_words;
     size_t mask;
 
-    // Hash functions ottimizzate - usa direttamente i 64 bytes interni
+    // Hash functions ottimizzate - usa direttamente i 64 bytes della chiave grezza X||Y
     inline uint64_t hash1(const uint8_t* data) const {
         const uint64_t* p = (const uint64_t*)data;
         return p[0] ^ (p[1] << 7);
@@ -117,8 +123,8 @@ public:
         free(bits);
     }
 
-    void add(const secp256k1_pubkey* pubkey) {
-        const uint8_t* data = pubkey->data;
+    // data: 64 byte grezzi X||Y (chiave pubblica non compressa senza prefisso 04)
+    void add(const uint8_t* data) {
         uint64_t h1 = hash1(data) & mask;
         uint64_t h2 = hash2(data) & mask;
         uint64_t h3 = hash3(data) & mask;
@@ -129,8 +135,7 @@ public:
     }
 
     // Verifica ultra-veloce con prefetching
-    inline bool might_contain(const secp256k1_pubkey* pubkey) const {
-        const uint8_t* data = pubkey->data;
+    inline bool might_contain(const uint8_t* data) const {
         uint64_t h1 = hash1(data) & mask;
         uint64_t h2 = hash2(data) & mask;
         uint64_t h3 = hash3(data) & mask;
@@ -156,7 +161,7 @@ static BloomFilter* bloom_filter = NULL;
 static volatile int keep_running = 1;
 static secp256k1_context* ctx = NULL;
 static std::vector<TargetKey> target_keys;
-static std::unordered_map<secp256k1_pubkey, int, PubkeyRawHash, PubkeyRawEqual> target_map;
+static std::unordered_map<std::array<uint8_t, 64>, int, RawKeyHash, RawKeyEqual> target_map;
 static uint64_t attempts_per_thread[MAX_THREADS] = {0};
 static time_t start_time;
 static pthread_mutex_t mutex = PTHREAD_MUTEX_INITIALIZER;
@@ -164,8 +169,13 @@ static FILE* log_file = NULL;
 static int num_threads = 0;
 
 #if USE_EC_BATCH
-// Precomputed: G, 2G, 3G, ..., 256G per batch EC addition
-static secp256k1_pubkey precomputed_G[EC_BATCH_SIZE];
+// Primo campo di secp256k1: p = 2^256 - 2^32 - 977
+static mpz_t field_p;
+
+// Multipli precalcolati del generatore: precomp_g[i] = (i+1)*G, per i = 0..EC_BATCH_MULT-1
+// (coordinate affini, condivisi in sola lettura tra i thread dopo l'inizializzazione)
+static mpz_t precomp_gx[EC_BATCH_MULT];
+static mpz_t precomp_gy[EC_BATCH_MULT];
 #endif
 
 // ============================================================================
@@ -273,15 +283,21 @@ int load_target_keys(const char* filename) {
         TargetKey key;
         if (hex_to_bytes(pubkey_hex.c_str(), key.pubkey, 65)) {
             strcpy(key.hex, pubkey_hex.c_str());
-            target_keys.push_back(key);
 
-            // Converti in secp256k1_pubkey per lookup diretto
+            // Valida che sia un punto valido sulla curva (scarta chiavi corrotte)
             secp256k1_pubkey pubkey_obj;
             if (secp256k1_ec_pubkey_parse(ctx, &pubkey_obj, key.pubkey, 65)) {
-                target_map[pubkey_obj] = count;
+                target_keys.push_back(key);
+
+                // Chiave di lookup: X||Y grezzi (64 byte), senza il prefisso 04.
+                // È lo stesso formato prodotto dal loop di generazione, quindi il
+                // confronto è un memcmp diretto, senza passare da secp256k1_pubkey.
+                std::array<uint8_t, 64> raw;
+                memcpy(raw.data(), key.pubkey + 1, 64);
+                target_map[raw] = count;
 
 #if USE_BLOOM_FILTER
-                bloom_filter->add(&pubkey_obj);
+                bloom_filter->add(raw.data());
 #endif
                 count++;
             }
@@ -294,29 +310,132 @@ int load_target_keys(const char* filename) {
     return count;
 }
 
+// ============================================================================
+// ARITMETICA DI CAMPO (GMP) PER BATCH EC POINT ADDITION
+// ============================================================================
+
+#if USE_EC_BATCH
+
+// Converte un intero mod p in 32 byte big-endian (con zero-padding)
+static inline void mpz_to_be32(const mpz_t x, uint8_t out[32]) {
+    memset(out, 0, 32);
+    size_t count = 0;
+    mpz_export(out, &count, 1, 1, 1, 0, x);
+    if (count > 0 && count < 32) {
+        memmove(out + (32 - count), out, count);
+        memset(out, 0, 32 - count);
+    }
+}
+
+// Converte 32 byte big-endian in un intero GMP
+static inline void be32_to_mpz(mpz_t x, const uint8_t in[32]) {
+    mpz_import(x, 32, 1, 1, 1, 0, in);
+}
+
+// Buffer di scratch riutilizzabili per l'addizione EC (evita alloc/dealloc per ogni chiave)
+struct EcScratch {
+    mpz_t H, HH, HHH, r, t1, t2;
+
+    void init() {
+        mpz_inits(H, HH, HHH, r, t1, t2, (mpz_ptr)0);
+    }
+    void clear_all() {
+        mpz_clears(H, HH, HHH, r, t1, t2, (mpz_ptr)0);
+    }
+};
+
+// Addizione di due punti affini (a=0, curva secp256k1), risultato in Jacobiane.
+// Formule "madd" specializzate per Z1=1 (il primo punto è sempre P0, appena
+// generato con una singola moltiplicazione scalare via libsecp256k1):
+//   H   = x2 - x1
+//   HH  = H^2
+//   HHH = H*HH
+//   r   = y2 - y1
+//   X3  = r^2 - HHH - 2*x1*HH
+//   Y3  = r*(x1*HH - X3) - y1*HHH
+//   Z3  = H
+//
+// NOTA: se H == 0 (x1 == x2 mod p, evento con probabilità ~2^-256 per punti
+// indipendenti) il risultato non è definito da queste formule; in quel caso
+// impostiamo Z3 = 0 per marcare il punto come "da ignorare" nel batch.
+static inline void ec_add_affine_affine(const mpz_t x1, const mpz_t y1,
+                                         const mpz_t x2, const mpz_t y2,
+                                         mpz_t X3, mpz_t Y3, mpz_t Z3,
+                                         EcScratch& s) {
+    mpz_sub(s.H, x2, x1);
+    mpz_mod(s.H, s.H, field_p);
+
+    if (mpz_sgn(s.H) == 0) {
+        mpz_set_ui(Z3, 0);
+        return;
+    }
+
+    mpz_mul(s.HH, s.H, s.H);
+    mpz_mod(s.HH, s.HH, field_p);
+
+    mpz_mul(s.HHH, s.H, s.HH);
+    mpz_mod(s.HHH, s.HHH, field_p);
+
+    mpz_sub(s.r, y2, y1);
+    mpz_mod(s.r, s.r, field_p);
+
+    // X3 = r^2 - HHH - 2*x1*HH
+    mpz_mul(X3, s.r, s.r);
+    mpz_sub(X3, X3, s.HHH);
+    mpz_mul(s.t1, x1, s.HH);
+    mpz_mul_2exp(s.t1, s.t1, 1);
+    mpz_sub(X3, X3, s.t1);
+    mpz_mod(X3, X3, field_p);
+
+    // Y3 = r*(x1*HH - X3) - y1*HHH
+    mpz_mul(s.t1, x1, s.HH);
+    mpz_sub(s.t1, s.t1, X3);
+    mpz_mul(Y3, s.r, s.t1);
+    mpz_mul(s.t2, y1, s.HHH);
+    mpz_sub(Y3, Y3, s.t2);
+    mpz_mod(Y3, Y3, field_p);
+
+    mpz_set(Z3, s.H);
+}
+
+// La batch inversion vera e propria (trucco di Montgomery: 1 sola mpz_invert
+// per l'intero batch, saltando gli eventuali punti non validi) è inline in
+// worker_thread, perché deve intrecciarsi con il vettore `valid[]`.
+
+#endif // USE_EC_BATCH
+
 // ============================================================================
 // PRECOMPUTE EC GENERATOR MULTIPLES
 // ============================================================================
 
 #if USE_EC_BATCH
 void precompute_generator_multiples() {
-    printf("[+] Precomputing EC generator multiples (1G, 2G, ..., %dG)...\n", EC_BATCH_SIZE);
+    printf("[+] Precomputing EC generator multiples (2G, ..., %dG)...\n", EC_BATCH_SIZE);
 
     uint8_t privkey[32];
 
-    for (int i = 0; i < EC_BATCH_SIZE; i++) {
+    for (int i = 0; i < EC_BATCH_MULT; i++) {
         memset(privkey, 0, 32);
 
-        // Imposta il valore (i+1) come privkey
-        // Per i=0: privkey=1, per i=255: privkey=256 (0x0100)
-        uint16_t value = i + 1;
-        privkey[31] = (uint8_t)(value & 0xFF);        // byte basso
-        privkey[30] = (uint8_t)((value >> 8) & 0xFF); // byte alto
+        // precomp_g[i] = (i+1)*G, per i=0..EC_BATCH_MULT-1 -> valori 1..EC_BATCH_MULT
+        uint32_t value = (uint32_t)(i + 1);
+        privkey[31] = (uint8_t)(value & 0xFF);
+        privkey[30] = (uint8_t)((value >> 8) & 0xFF);
 
-        if (!secp256k1_ec_pubkey_create(ctx, &precomputed_G[i], privkey)) {
-            fprintf(stderr, "[ERROR] Failed to precompute %dG\n", i+1);
+        secp256k1_pubkey pk;
+        if (!secp256k1_ec_pubkey_create(ctx, &pk, privkey)) {
+            fprintf(stderr, "[ERROR] Failed to precompute %dG\n", i + 1);
             exit(1);
         }
+
+        unsigned char buf[65];
+        size_t outlen = 65;
+        secp256k1_ec_pubkey_serialize(ctx, buf, &outlen, &pk, SECP256K1_EC_UNCOMPRESSED);
+
+        mpz_init(precomp_gx[i]);
+        mpz_init(precomp_gy[i]);
+        be32_to_mpz(precomp_gx[i], buf + 1);
+        be32_to_mpz(precomp_gy[i], buf + 33);
     }
 
     printf("[+] Precomputation complete!\n");
@@ -368,20 +487,40 @@ static inline void add_to_privkey(uint8_t* privkey, uint64_t n) {
     }
 }
 
+// Somma un piccolo intero a un array di 32 byte interpretato come intero
+// big-endian standard (byte[31] = LSB), lo stesso formato che secp256k1 si
+// aspetta per una private key. NON intercambiabile con add_to_privkey/
+// increment_privkey sopra: quelle due, per velocità, trattano l'array come
+// 4 word a 64 bit in ordine nativo (little-endian sulle CPU x86), quindi
+// avanzano la ricerca in un ordine "rimescolato" ma comunque bigettivo sullo
+// spazio delle chiavi — va benissimo per iterare, ma NON per ricostruire lo
+// scalare esatto corrispondente a un punto EC calcolato con aritmetica reale
+// (come batch[i] = P0 + (i+1)*G). Per quello serve questa versione corretta.
+static inline void add_small_be256(uint8_t* be, uint32_t n) {
+    for (int i = 31; i >= 0 && n; i--) {
+        uint32_t sum = be[i] + (n & 0xFF);
+        be[i] = (uint8_t)sum;
+        n = (n >> 8) + (sum >> 8);
+    }
+}
+
 // ============================================================================
 // MATCH CHECKING OTTIMIZZATO
 // ============================================================================
 
-static inline int check_match_fast(const secp256k1_pubkey* pubkey) {
+// data: 64 byte grezzi X||Y (chiave pubblica non compressa senza prefisso 04)
+static inline int check_match_fast_raw(const uint8_t* data) {
 #if USE_BLOOM_FILTER
     // Prima passa: Bloom filter
-    if (!bloom_filter->might_contain(pubkey)) {
+    if (!bloom_filter->might_contain(data)) {
         return -1; // Sicuramente non presente
     }
 #endif
 
-    // Lookup diretto nella hash map (zero copy!)
-    auto it = target_map.find(*pubkey);
+    std::array<uint8_t, 64> key;
+    memcpy(key.data(), data, 64);
+
+    auto it = target_map.find(key);
     if (it != target_map.end()) {
         return it->second; // Indice nella lista target_keys
     }
@@ -438,33 +577,56 @@ void format_number(uint64_t num, char* buffer) {
     }
 }
 
+// Velocità istantanea: somma dei tentativi di tutti i thread nella finestra
+// trascorsa dall'ultima chiamata (non media cumulata dall'avvio), così il
+// numero mostrato riflette il ritmo REALE corrente anche su run lunghe.
 void log_progress() {
     pthread_mutex_lock(&mutex);
 
+    static uint64_t last_total = 0;
+    static struct timeval last_tv = {0, 0};
+
+    struct timeval now_tv;
+    gettimeofday(&now_tv, NULL);
+    if (last_tv.tv_sec == 0 && last_tv.tv_usec == 0) {
+        // Prima chiamata: usa l'avvio dei thread come inizio finestra, non "adesso"
+        // (altrimenti la finestra sarebbe ~0s e il rate esploderebbe verso l'infinito)
+        last_tv.tv_sec = start_time;
+        last_tv.tv_usec = 0;
+    }
+
     time_t now = time(NULL);
-    double elapsed = difftime(now, start_time);
-    if (elapsed < 1) elapsed = 1;
+    double elapsed_total = difftime(now, start_time);
+    if (elapsed_total < 1) elapsed_total = 1;
 
     uint64_t total = 0;
     for (int i = 0; i < num_threads; i++) {
         total += attempts_per_thread[i];
     }
 
-    double rate = total / elapsed;
+    double window_sec = (now_tv.tv_sec - last_tv.tv_sec) +
+                         (now_tv.tv_usec - last_tv.tv_usec) / 1e6;
+    if (window_sec < 0.001) window_sec = 0.001;
+
+    uint64_t delta = total - last_total;
+    double instant_rate = delta / window_sec;
 
     char total_str[32];
     char rate_str[32];
     format_number(total, total_str);
-    format_number((uint64_t)rate, rate_str);
+    format_number((uint64_t)instant_rate, rate_str);
 
-    printf("[INFO] Tentativi: %s | Velocità: %s keys/sec | Tempo: %.0fs\n",
-           total_str, rate_str, elapsed);
+    printf("[INFO] Tentativi totali: %s | Velocità: %s keys/sec (tutti i core) | Tempo: %.0fs\n",
+           total_str, rate_str, elapsed_total);
 
     if (log_file) {
-        fprintf(log_file, "%ld,%lu,%.2f\n", now, total, rate);
+        fprintf(log_file, "%ld,%lu,%.2f\n", now, total, instant_rate);
         fflush(log_file);
     }
 
+    last_total = total;
+    last_tv = now_tv;
+
     pthread_mutex_unlock(&mutex);
 }
 
@@ -481,7 +643,6 @@ void* worker_thread(void* arg) {
 
     // Pre-alloca buffer
     uint8_t privkey[32];
-    secp256k1_pubkey pubkey_batch[EC_BATCH_SIZE];
     uint64_t local_attempts = 0;
 
     init_random_privkey_in_range(privkey, &seed, data->range_start, data->range_end);
@@ -492,41 +653,113 @@ void* worker_thread(void* arg) {
     printf("    Privkey iniziale: %s\n", privkey_start_hex);
 
     // ========================================================================
-    // LOOP PRINCIPALE CON EC BATCH PROCESSING
+    // LOOP PRINCIPALE CON EC BATCH PROCESSING (Jacobian + batch inversion)
     // ========================================================================
 
 #if USE_EC_BATCH
-    // VERSIONE CON BATCH EC POINT ADDITION
+    // Buffer persistenti per thread: X/Y/Z Jacobiani dei punti P0 + iG (i=1..EC_BATCH_MULT)
+    mpz_t x0, y0;
+    mpz_t Xj[EC_BATCH_MULT], Yj[EC_BATCH_MULT], Zj[EC_BATCH_MULT];
+    mpz_t invZ[EC_BATCH_MULT], prefix[EC_BATCH_MULT];
+    mpz_t inv_tmp, zinv2, zinv3, ax, ay;
+    EcScratch scratch;
+
+    mpz_inits(x0, y0, inv_tmp, zinv2, zinv3, ax, ay, (mpz_ptr)0);
+    for (int i = 0; i < EC_BATCH_MULT; i++) {
+        mpz_inits(Xj[i], Yj[i], Zj[i], invZ[i], prefix[i], (mpz_ptr)0);
+    }
+    scratch.init();
+
+    uint8_t rawkey[64];
+    uint8_t found_privkey[32];
+
     while (keep_running) {
-        // Step 1: Genera la prima pubkey del batch (P = privkey * G)
-        if (!secp256k1_ec_pubkey_create(ctx, &pubkey_batch[0], privkey)) {
+        secp256k1_pubkey pubkey0;
+        if (!secp256k1_ec_pubkey_create(ctx, &pubkey0, privkey)) {
             increment_privkey(privkey);
             continue;
         }
 
-        // Step 2: Check prima chiave
-        int match_idx = check_match_fast(&pubkey_batch[0]);
+        unsigned char buf65[65];
+        size_t outlen = 65;
+        secp256k1_ec_pubkey_serialize(ctx, buf65, &outlen, &pubkey0, SECP256K1_EC_UNCOMPRESSED);
+
+        be32_to_mpz(x0, buf65 + 1);
+        be32_to_mpz(y0, buf65 + 33);
+
+        // Chiave 0 del batch: è P0 stesso, nessuna conversione affine necessaria
+        int match_idx = check_match_fast_raw(buf65 + 1);
         if (__builtin_expect(match_idx >= 0, 0)) {
             save_found_key(privkey, match_idx);
         }
 
-        // Step 3: Genera le restanti (EC_BATCH_SIZE - 1) chiavi usando EC addition
-        // P1 = P + G, P2 = P + 2G, P3 = P + 3G, ...
-        // Questo è MOLTO più veloce di fare EC_BATCH_SIZE moltiplicazioni!
-        uint8_t temp_privkey[32];
-        memcpy(temp_privkey, privkey, 32);
+        // Genera le restanti EC_BATCH_MULT chiavi: batch[i] = P0 + (i+1)*G, in Jacobiane
+        int valid_count = 0;
+        bool valid[EC_BATCH_MULT];
+        for (int i = 0; i < EC_BATCH_MULT; i++) {
+            ec_add_affine_affine(x0, y0, precomp_gx[i], precomp_gy[i],
+                                  Xj[i], Yj[i], Zj[i], scratch);
+            valid[i] = (mpz_sgn(Zj[i]) != 0);
+            if (valid[i]) valid_count++;
+        }
 
-        for (int i = 1; i < EC_BATCH_SIZE && keep_running; i++) {
-            increment_privkey(temp_privkey);
+        // Batch inversion (trucco di Montgomery): una sola mpz_invert per l'intero batch
+        if (valid_count > 0) {
+            int last = -1;
+            for (int i = 0; i < EC_BATCH_MULT; i++) {
+                if (!valid[i]) continue;
+                if (last < 0) {
+                    mpz_set(prefix[i], Zj[i]);
+                } else {
+                    mpz_mul(prefix[i], prefix[last], Zj[i]);
+                    mpz_mod(prefix[i], prefix[i], field_p);
+                }
+                last = i;
+            }
 
-            // EC point addition: pubkey_batch[i] = pubkey_batch[0] + precomputed_G[i-1]
-            // Usa EC pubkey combine (somma di due punti)
-            const secp256k1_pubkey* pubkeys_to_add[2] = {&pubkey_batch[0], &precomputed_G[i]};
+            mpz_invert(inv_tmp, prefix[last], field_p); // inv_tmp = 1 / prodotto totale
 
-            if (secp256k1_ec_pubkey_combine(ctx, &pubkey_batch[i], pubkeys_to_add, 2)) {
-                match_idx = check_match_fast(&pubkey_batch[i]);
+            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; }
+                }
+
+                if (prev >= 0) {
+                    mpz_mul(invZ[i], inv_tmp, prefix[prev]);
+                    mpz_mod(invZ[i], invZ[i], field_p);
+                } else {
+                    mpz_set(invZ[i], inv_tmp);
+                }
+
+                mpz_mul(inv_tmp, inv_tmp, Zj[i]);
+                mpz_mod(inv_tmp, inv_tmp, field_p);
+            }
+
+            // Converte ogni punto Jacobiano in affine e verifica il match
+            for (int i = 0; i < EC_BATCH_MULT && keep_running; i++) {
+                if (!valid[i]) continue;
+
+                mpz_mul(zinv2, invZ[i], invZ[i]);
+                mpz_mod(zinv2, zinv2, field_p);
+                mpz_mul(zinv3, zinv2, invZ[i]);
+                mpz_mod(zinv3, zinv3, field_p);
+
+                mpz_mul(ax, Xj[i], zinv2);
+                mpz_mod(ax, ax, field_p);
+                mpz_mul(ay, Yj[i], zinv3);
+                mpz_mod(ay, ay, field_p);
+
+                mpz_to_be32(ax, rawkey);
+                mpz_to_be32(ay, rawkey + 32);
+
+                match_idx = check_match_fast_raw(rawkey);
                 if (__builtin_expect(match_idx >= 0, 0)) {
-                    save_found_key(temp_privkey, match_idx);
+                    memcpy(found_privkey, privkey, 32);
+                    add_small_be256(found_privkey, (uint32_t)(i + 1)); // batch[i] = privkey + (i+1)
+                    save_found_key(found_privkey, match_idx);
                 }
             }
         }
@@ -535,10 +768,17 @@ void* worker_thread(void* arg) {
         add_to_privkey(privkey, EC_BATCH_SIZE);
 
         // Aggiorna contatore globale periodicamente
-        if ((local_attempts & (SYNC_BATCH - 1)) == 0) {
+        // (modulo, non maschera bitwise: SYNC_BATCH non è una potenza di due)
+        if (local_attempts % SYNC_BATCH < EC_BATCH_SIZE) {
             attempts_per_thread[thread_id] = local_attempts;
         }
     }
+
+    scratch.clear_all();
+    mpz_clears(x0, y0, inv_tmp, zinv2, zinv3, ax, ay, (mpz_ptr)0);
+    for (int i = 0; i < EC_BATCH_MULT; i++) {
+        mpz_clears(Xj[i], Yj[i], Zj[i], invZ[i], prefix[i], (mpz_ptr)0);
+    }
 #else
     // VERSIONE STANDARD (fallback senza batch)
     while (keep_running) {
@@ -546,7 +786,10 @@ void* worker_thread(void* arg) {
 
         for (int batch = 0; batch < SYNC_BATCH; batch++) {
             if (__builtin_expect(secp256k1_ec_pubkey_create(ctx, &pubkey_obj, privkey), 1)) {
-                int match_idx = check_match_fast(&pubkey_obj);
+                unsigned char buf65[65];
+                size_t outlen = 65;
+                secp256k1_ec_pubkey_serialize(ctx, buf65, &outlen, &pubkey_obj, SECP256K1_EC_UNCOMPRESSED);
+                int match_idx = check_match_fast_raw(buf65 + 1);
                 if (__builtin_expect(match_idx >= 0, 0)) {
                     save_found_key(privkey, match_idx);
                 }
@@ -602,8 +845,12 @@ int main(int argc, char** argv) {
         }
     }
 
-    // Precompute EC multiples
 #if USE_EC_BATCH
+    // Inizializza il primo campo di secp256k1: p = 2^256 - 2^32 - 977
+    mpz_init_set_str(field_p,
+        "FFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFEFFFFFC2F", 16);
+
+    // Precompute EC multiples
     precompute_generator_multiples();
 #endif
 
@@ -662,7 +909,7 @@ int main(int argc, char** argv) {
 
     // Loop principale
     while (keep_running) {
-        sleep(10);
+        sleep(PROGRESS_INTERVAL_SEC);
         log_progress();
     }
 

bruteforce/p2pk_bruteforce_gpu.cu

@@ -0,0 +1,727 @@
+/*
+ * 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;
+}