From ac68b6aa883000e95d6b20abfdc8a73ee4c9f51b Mon Sep 17 00:00:00 2001 From: Davide Grilli Date: Fri, 3 Jul 2026 15:58:26 +0200 Subject: [PATCH] Aggiungi versione GPU (CUDA) del bruteforce MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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 | 1 + CLAUDE.md | 24 +- bruteforce/Makefile | 50 +- bruteforce/README.md | 26 ++ bruteforce/p2pk_bruteforce_gpu.cu | 727 ++++++++++++++++++++++++++++++ 5 files changed, 820 insertions(+), 8 deletions(-) create mode 100644 bruteforce/p2pk_bruteforce_gpu.cu diff --git a/.gitignore b/.gitignore index 76b3402..606871e 100644 --- a/.gitignore +++ b/.gitignore @@ -29,6 +29,7 @@ build/ # Eseguibili compilati bruteforce/p2pk_bruteforce bruteforce/p2pk_bruteforce_debug +bruteforce/p2pk_bruteforce_gpu # File oggetto e compilazione *.o diff --git a/CLAUDE.md b/CLAUDE.md index b873cf9..5b668d2 100644 --- a/CLAUDE.md +++ b/CLAUDE.md @@ -7,7 +7,7 @@ This file provides guidance to Claude Code (claude.ai/code) when working with co 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/` — C++ program that loads target public keys and searches the private-key space for matches, as a performance demonstration (not a realistic attack — success probability is ~2^-256). +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. @@ -41,18 +41,32 @@ python3 extract_p2pk_utxo.py [db_path] [output.txt] # default: ../databases/bi 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` (single file, everything lives here) +### `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` (exact match) and a 64MB Bloom filter (fast negative rejection, checked first via `check_match_fast`). -- Performance trick: instead of computing `privkey * G` per candidate, it computes one EC multiplication then derives the next `EC_BATCH_SIZE` (256) keys via EC point addition against precomputed multiples of G (`precompute_generator_multiples`), which is much cheaper than repeated scalar multiplication. +- Target pubkeys are loaded from a `.txt` file (uncompressed hex, `04` prefix, one per line, header line skipped) into both an `unordered_map,...>` 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`; periodic throughput stats go to stdout and `progress.csv` every 10s. +- 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. diff --git a/bruteforce/Makefile b/bruteforce/Makefile index 5a24cb9..94d0939 100644 --- a/bruteforce/Makefile +++ b/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 diff --git a/bruteforce/README.md b/bruteforce/README.md index bff7e26..5215d26 100644 --- a/bruteforce/README.md +++ b/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 diff --git a/bruteforce/p2pk_bruteforce_gpu.cu b/bruteforce/p2pk_bruteforce_gpu.cu new file mode 100644 index 0000000..5baec18 --- /dev/null +++ b/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 +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +// ============================================================================ +// 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 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 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 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 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<<>>( + 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; +}