From 6bee6b965e93c18871ff7fcf105c0234a416e6d3 Mon Sep 17 00:00:00 2001 From: cavey Date: Thu, 10 Sep 2026 21:26:18 +0000 Subject: [PATCH 1/2] add --pattern ORs using separate one-prefix and multi CUDA kernels --- README.md | 33 ++-- build.rs | 4 +- kernels/base58.cu | 25 +++ kernels/base58.h | 115 ++++++++++++ kernels/opencl/base58.cl | 106 ++++++++++- kernels/opencl/host.c | 96 +++------- kernels/opencl/keypair.cl | 22 ++- kernels/opencl/vanity.cl | 22 ++- kernels/pattern_table.h | 42 +++++ kernels/vanity.cu | 89 +++++---- kernels/vanity.h | 4 +- kernels/vanity_keypair.cu | 105 ++++++----- kernels/vanity_keypair.h | 2 +- src/fast/check_match.rs | 225 +++++++++++++++-------- src/fast/mod.rs | 16 +- src/main.rs | 376 +++++++++++++++++++++++++++----------- 16 files changed, 913 insertions(+), 369 deletions(-) create mode 100644 kernels/pattern_table.h diff --git a/README.md b/README.md index e4081eb..d04321d 100644 --- a/README.md +++ b/README.md @@ -59,13 +59,23 @@ No ed25519 is involved, so it's dramatically faster — but the account is controlled by `base` (the signer), not by a standalone keypair. You supply the `--base` pubkey (the signer for the create instruction), the -program `--owner`, and a `--prefix` and/or `--suffix` to match (base58): +program `--owner`, and one or more `--pattern`s to match (base58). A pubkey +hits if it matches **any** pattern. `...` splits prefix and suffix (`.` is +not in base58): ```bash vanity grind \ --base \ --owner BPFLoaderUpgradeab1e11111111111111111111111 \ - --prefix abc \ + --pattern abc... \ + --num-gpus 1 + +# several targets in one run (OR): +vanity grind \ + --base \ + --owner BPFLoaderUpgradeab1e11111111111111111111111 \ + --pattern Cavey...CooL \ + --pattern Harmonic... \ --num-gpus 1 ``` @@ -103,15 +113,16 @@ pub fn create_account_with_seed( ### Grind a vanity keypair -`grind-keypair` searches genuine ed25519 keypairs whose **public key** matches a -base58 `--prefix` and/or `--suffix`. This is the classic "vanity wallet/mint" -case: the result is a standalone keypair you can sign with. It's slower than -`grind` because each attempt computes an ed25519 scalar multiplication. The CPU -path uses a batched custom ed25519 implementation (AVX-512 IFMA when available); -GPUs use the CUDA/OpenCL keypair kernels. +`grind-keypair` searches genuine ed25519 keypairs whose **public key** matches +any `--pattern` (same `Prefix...Suffix` syntax as `grind`). This is the classic +"vanity wallet/mint" case: the result is a standalone keypair you can sign +with. It's slower than `grind` because each attempt computes an ed25519 scalar +multiplication. The CPU path uses a batched custom ed25519 implementation +(AVX-512 IFMA when available); GPUs use the CUDA/OpenCL keypair kernels. ```bash -vanity grind-keypair --prefix bob --suffix xyz --num-gpus 1 +vanity grind-keypair --pattern bob...xyz --num-gpus 1 +vanity grind-keypair --pattern Cavey...CooL --pattern Harmonic... --num-gpus 1 ``` On a match it prints the public key, the seed (hex), and a Solana-compatible @@ -176,8 +187,8 @@ doppler: 1/4 sign-extendable segment(s) | `--num-cpus ` | all | `0` (= all logical cores) | CPU mining threads | | `--num-gpus ` | all (GPU builds only) | `1` | GPUs to mine on | | `--count ` | all | `1` | stop after finding N matches | -| `--case-insensitive` | `grind`, `grind-keypair` | off | match prefix/suffix ignoring case | -| `--prefix` / `--suffix` | `grind`, `grind-keypair` | — | base58 target(s) to match; supply at least one | +| `--case-insensitive` | `grind`, `grind-keypair` | off | match patterns ignoring case | +| `--pattern` | `grind`, `grind-keypair` | — | repeatable `Prefix...Suffix` / `Prefix...` / `...Suffix`; OR'd | | `--segments <1-4>` | `grind-doppler` | `1` | sign-extendable segments required | To run purely on CPU (no GPU), build without a GPU feature, or pass `--num-gpus 0`. diff --git a/build.rs b/build.rs index 83f0ed3..bf52293 100644 --- a/build.rs +++ b/build.rs @@ -86,6 +86,7 @@ fn build_opencl_libs() { use std::path::Path; println!("cargo::rerun-if-changed=kernels/opencl/"); + println!("cargo::rerun-if-changed=kernels/pattern_table.h"); let out_dir = std::env::var("OUT_DIR").unwrap(); let target_os = @@ -130,7 +131,8 @@ fn build_opencl_libs() { let mut build = cc::Build::new(); build .file("kernels/opencl/host.c") - .include(&out_dir); + .include(&out_dir) + .include("kernels"); if target_os == "macos" { build.flag("-Wno-deprecated-declarations"); // Pin the SDK so the OpenCL.framework headers resolve regardless of diff --git a/kernels/base58.cu b/kernels/base58.cu index b38d453..becf975 100644 --- a/kernels/base58.cu +++ b/kernels/base58.cu @@ -1,3 +1,5 @@ +#include +#include #include "base58.h" __device__ uint8_t const base58_chars[] = "123456789ABCDEFGHJKLMNPQRSTUVWXYZabcdefghijkmnopqrstuvwxyz"; @@ -11,6 +13,29 @@ __device__ uint8_t const base58_chars_ci[] = "123456789abcdefghjkLmnpqrstuvwxyza is always raw_base58[i] -> d_match_lut[...] vs precomputed target. */ __constant__ uint8_t d_match_lut[58]; +#include "pattern_table.h" + +__constant__ uint32_t d_n_patterns; +__constant__ uint8_t d_prefix_len[VANITY_MAX_PATTERNS]; +__constant__ uint8_t d_suffix_len[VANITY_MAX_PATTERNS]; +__constant__ uint8_t d_prefix[VANITY_MAX_PATTERNS][VANITY_MAX_PATTERN_LEN]; +__constant__ uint8_t d_suffix[VANITY_MAX_PATTERNS][VANITY_MAX_PATTERN_LEN]; + +extern "C" void vanity_upload_pattern_table(const uint8_t *blob, uint64_t blob_len) +{ + uint8_t flat[VANITY_PT_SIZE]; + vanity_unpack_patterns(blob, blob_len, flat); + uint32_t n = 0; + memcpy(&n, flat + VANITY_PT_N, 4); + cudaMemcpyToSymbol(d_n_patterns, &n, 4); + cudaMemcpyToSymbol(d_prefix_len, flat + VANITY_PT_PLEN, VANITY_MAX_PATTERNS); + cudaMemcpyToSymbol(d_suffix_len, flat + VANITY_PT_SLEN, VANITY_MAX_PATTERNS); + cudaMemcpyToSymbol(d_prefix, flat + VANITY_PT_PREF, + VANITY_MAX_PATTERNS * VANITY_MAX_PATTERN_LEN); + cudaMemcpyToSymbol(d_suffix, flat + VANITY_PT_SUF, + VANITY_MAX_PATTERNS * VANITY_MAX_PATTERN_LEN); +} + #define BASE58_INVALID_CHAR ((uint8_t)255) #define BASE58_INVERSE_TABLE_OFFSET ((uint8_t)'1') #define BASE58_INVERSE_TABLE_SENTINEL ((uint8_t)(1UL + (uint8_t)('z') - BASE58_INVERSE_TABLE_OFFSET)) diff --git a/kernels/base58.h b/kernels/base58.h index 6695df7..4ea6bc0 100644 --- a/kernels/base58.h +++ b/kernels/base58.h @@ -10,6 +10,22 @@ extern __device__ uint8_t const base58_chars_ci[]; extern __device__ uint const enc_table_32[8][8]; extern __constant__ uint8_t d_match_lut[58]; +#ifndef VANITY_MAX_PATTERNS +#define VANITY_MAX_PATTERNS 64 +#define VANITY_MAX_PATTERN_LEN 44 +#endif + +extern __constant__ uint32_t d_n_patterns; +extern __constant__ uint8_t d_prefix_len[VANITY_MAX_PATTERNS]; +extern __constant__ uint8_t d_suffix_len[VANITY_MAX_PATTERNS]; +extern __constant__ uint8_t d_prefix[VANITY_MAX_PATTERNS][VANITY_MAX_PATTERN_LEN]; +extern __constant__ uint8_t d_suffix[VANITY_MAX_PATTERNS][VANITY_MAX_PATTERN_LEN]; + +#ifdef __cplusplus +extern "C" +#endif +void vanity_upload_pattern_table(const uint8_t *blob, uint64_t blob_len); + __device__ ulong fd_base58_encode_32(uint8_t *bytes, uint8_t *out, bool case_insensitive); /* Word-form fused base58 encode + prefix/suffix match with early @@ -112,4 +128,103 @@ static __device__ __forceinline__ bool fd_base58_check_match_32_words(const uint #undef VANITY_BS58_ENSURE_LIMB } +/* Multi-pattern: encode once, OR each pair from __constant__ SoA. + Forceinlined only into the MULTI=true kernel (the 1-pattern kernel + never calls this), so its registers cannot leak into that occupancy. */ +static __device__ __forceinline__ bool fd_base58_check_match_any_32_words(const uint state[8]) +{ + ulong in_leading_0s = 0UL; +#pragma unroll + for (int i = 0; i < 8; ++i) + { + ulong lz_bytes = (ulong)((unsigned)__clz((int)state[i]) >> 3); + in_leading_0s += lz_bytes; + if (lz_bytes != 4UL) break; + } + + ulong R1div = 656356768UL; + ulong intermediate[9] = {0}; + for (ulong i = 0UL; i < 8UL; i++) + for (ulong j = 0UL; j < 8UL; j++) + intermediate[j + 1UL] += (ulong)state[i] * (ulong)enc_table_32[i][j]; + + for (ulong i = 8UL; i > 0UL; i--) + { + intermediate[i - 1UL] += (intermediate[i] / R1div); + intermediate[i] %= R1div; + } + + uint8_t raw_base58[45]; + ulong limbs_done = 0UL; + #define VANITY_BS58_EMIT_LIMB_M(L) do { \ + uint v = (uint)intermediate[(L)]; \ + raw_base58[5UL*(L) + 4UL] = (uint8_t)((v / 1U) % 58U); \ + raw_base58[5UL*(L) + 3UL] = (uint8_t)((v / 58U) % 58U); \ + raw_base58[5UL*(L) + 2UL] = (uint8_t)((v / 3364U) % 58U); \ + raw_base58[5UL*(L) + 1UL] = (uint8_t)((v / 195112U) % 58U); \ + raw_base58[5UL*(L) + 0UL] = (uint8_t)(v / 11316496U); \ + } while (0) + #define VANITY_BS58_ENSURE_LIMB_M(L) do { \ + while (limbs_done <= (L)) { VANITY_BS58_EMIT_LIMB_M(limbs_done); limbs_done++; } \ + } while (0) + + VANITY_BS58_ENSURE_LIMB_M(0UL); + + ulong raw_leading_0s = 0UL; + while (raw_leading_0s < 45UL) + { + if (raw_leading_0s / 5UL >= limbs_done) + VANITY_BS58_ENSURE_LIMB_M(raw_leading_0s / 5UL); + if (raw_base58[raw_leading_0s]) + break; + raw_leading_0s++; + } + + ulong skip = raw_leading_0s - in_leading_0s; + ulong encoded_length = 45UL - skip; + + uint32_t n_pat = d_n_patterns; + if (n_pat == 0) + return true; + + for (uint32_t pi = 0; pi < n_pat; pi++) + { + uint8_t plen = d_prefix_len[pi]; + uint8_t slen = d_suffix_len[pi]; + + bool ok = true; + for (ulong i = 0UL; i < (ulong)plen; i++) + { + ulong rb_idx = skip + i; + VANITY_BS58_ENSURE_LIMB_M(rb_idx / 5UL); + if (d_match_lut[raw_base58[rb_idx]] != d_prefix[pi][i]) + { + ok = false; + break; + } + } + if (ok && slen > 0) + { + ulong tail_start = skip + encoded_length - (ulong)slen; + ulong last_limb = (skip + encoded_length - 1UL) / 5UL; + VANITY_BS58_ENSURE_LIMB_M(last_limb); + for (ulong i = 0UL; i < (ulong)slen; i++) + { + if (d_match_lut[raw_base58[tail_start + i]] != d_suffix[pi][i]) + { + ok = false; + break; + } + } + } + if (ok) + return true; + } + + return false; + + #undef VANITY_BS58_EMIT_LIMB_M + #undef VANITY_BS58_ENSURE_LIMB_M +} + #endif diff --git a/kernels/opencl/base58.cl b/kernels/opencl/base58.cl index 3160fca..af7241e 100644 --- a/kernels/opencl/base58.cl +++ b/kernels/opencl/base58.cl @@ -6,6 +6,16 @@ word form, with early rejection (grind + keypair kernels). Both descend from Firedancer's fd_base58 (cavemanloverboy port). */ +#ifndef VANITY_MAX_PATTERNS +#define VANITY_MAX_PATTERNS 64 +#define VANITY_MAX_PATTERN_LEN 44 +#define VANITY_PT_N 0 +#define VANITY_PT_PLEN 4 +#define VANITY_PT_SLEN (4 + VANITY_MAX_PATTERNS) +#define VANITY_PT_PREF (VANITY_PT_SLEN + VANITY_MAX_PATTERNS) +#define VANITY_PT_SUF (VANITY_PT_PREF + VANITY_MAX_PATTERNS * VANITY_MAX_PATTERN_LEN) +#endif + __constant uchar base58_chars[] = "123456789ABCDEFGHJKLMNPQRSTUVWXYZabcdefghijkmnopqrstuvwxyz"; /* ci = case insensitive: map all uppercase to lowercase except for L. */ __constant uchar base58_chars_ci[] = "123456789abcdefghjkLmnpqrstuvwxyzabcdefghijkmnopqrstuvwxyz"; @@ -79,8 +89,8 @@ static ulong fd_base58_encode_32(const uchar *bytes, uchar *out, bool case_insen of canonical raw_base58 indices precomputed on the host; match_lut folds raw_base58 values into the same canonical space. */ static bool fd_base58_check_match_32_words(const uint state[8], - __global const uchar *target, ulong target_len, - __global const uchar *suffix, ulong suffix_len, + __constant const uchar *target, ulong target_len, + __constant const uchar *suffix, ulong suffix_len, __constant const uchar *match_lut) { /* Count leading zero bytes of the big-endian byte view of state[]. */ ulong in_leading_0s = 0UL; @@ -151,3 +161,95 @@ static bool fd_base58_check_match_32_words(const uint state[8], #undef VANITY_BS58_EMIT_LIMB #undef VANITY_BS58_ENSURE_LIMB } + +/* Multi-pattern: encode once, OR each packed (prefix, suffix) pair. */ +static bool fd_base58_check_match_any_32_words(const uint state[8], + __constant const uchar *patterns, + __constant const uchar *match_lut) { + ulong in_leading_0s = 0UL; + for (int i = 0; i < 8; ++i) { + ulong lz_bytes = (ulong)(clz(state[i]) >> 3); + in_leading_0s += lz_bytes; + if (lz_bytes != 4UL) break; + } + + ulong R1div = 656356768UL; + ulong intermediate[9] = {0}; + for (ulong i = 0UL; i < 8UL; i++) + for (ulong j = 0UL; j < 8UL; j++) + intermediate[j + 1UL] += (ulong)state[i] * (ulong)enc_table_32[i][j]; + + for (ulong i = 8UL; i > 0UL; i--) { + intermediate[i - 1UL] += (intermediate[i] / R1div); + intermediate[i] %= R1div; + } + + uchar raw_base58[45]; + ulong limbs_done = 0UL; + #define VANITY_BS58_EMIT_LIMB_M(L) do { \ + uint v = (uint)intermediate[(L)]; \ + raw_base58[5UL*(L) + 4UL] = (uchar)((v / 1U) % 58U); \ + raw_base58[5UL*(L) + 3UL] = (uchar)((v / 58U) % 58U); \ + raw_base58[5UL*(L) + 2UL] = (uchar)((v / 3364U) % 58U); \ + raw_base58[5UL*(L) + 1UL] = (uchar)((v / 195112U) % 58U); \ + raw_base58[5UL*(L) + 0UL] = (uchar)( v / 11316496U); \ + } while (0) + #define VANITY_BS58_ENSURE_LIMB_M(L) do { \ + while (limbs_done <= (L)) { VANITY_BS58_EMIT_LIMB_M(limbs_done); limbs_done++; } \ + } while (0) + + VANITY_BS58_ENSURE_LIMB_M(0UL); + + ulong raw_leading_0s = 0UL; + while (raw_leading_0s < 45UL) { + if (raw_leading_0s / 5UL >= limbs_done) + VANITY_BS58_ENSURE_LIMB_M(raw_leading_0s / 5UL); + if (raw_base58[raw_leading_0s]) break; + raw_leading_0s++; + } + + ulong skip = raw_leading_0s - in_leading_0s; + ulong encoded_length = 45UL - skip; + + uint n_pat = (uint)patterns[VANITY_PT_N] + | ((uint)patterns[VANITY_PT_N + 1] << 8) + | ((uint)patterns[VANITY_PT_N + 2] << 16) + | ((uint)patterns[VANITY_PT_N + 3] << 24); + if (n_pat == 0) + return true; + + for (uint pi = 0; pi < n_pat; pi++) { + uchar plen = patterns[VANITY_PT_PLEN + pi]; + uchar slen = patterns[VANITY_PT_SLEN + pi]; + ulong pref = (ulong)VANITY_PT_PREF + (ulong)pi * (ulong)VANITY_MAX_PATTERN_LEN; + ulong suf = (ulong)VANITY_PT_SUF + (ulong)pi * (ulong)VANITY_MAX_PATTERN_LEN; + + bool ok = true; + for (ulong i = 0UL; i < (ulong)plen; i++) { + ulong rb_idx = skip + i; + VANITY_BS58_ENSURE_LIMB_M(rb_idx / 5UL); + if (match_lut[raw_base58[rb_idx]] != patterns[pref + i]) { + ok = false; + break; + } + } + if (ok && slen > 0) { + ulong tail_start = skip + encoded_length - (ulong)slen; + ulong last_limb = (skip + encoded_length - 1UL) / 5UL; + VANITY_BS58_ENSURE_LIMB_M(last_limb); + for (ulong i = 0UL; i < (ulong)slen; i++) { + if (match_lut[raw_base58[tail_start + i]] != patterns[suf + i]) { + ok = false; + break; + } + } + } + if (ok) + return true; + } + + return false; + + #undef VANITY_BS58_EMIT_LIMB_M + #undef VANITY_BS58_ENSURE_LIMB_M +} diff --git a/kernels/opencl/host.c b/kernels/opencl/host.c index 36c701b..b2b9aca 100644 --- a/kernels/opencl/host.c +++ b/kernels/opencl/host.c @@ -27,6 +27,7 @@ #include #include "cl_sources.h" +#include "pattern_table.h" /* ─── tunables ──────────────────────────────────────────────────────────── */ @@ -176,9 +177,8 @@ typedef struct { cl_command_queue queue; cl_program program; cl_kernel kernel; - cl_mem seed, w0, sr7, w1, glyph, mlut, target, suffix, out, done, counts; + cl_mem seed, w0, sr7, w1, glyph, mlut, patterns, out, done, counts; size_t local, global; - uint32_t target_len, suffix_len; uint32_t max_iters; cl_event event; int in_flight; @@ -187,8 +187,7 @@ typedef struct { } GrindCtx; void *gpu_grind_init(int id, uint8_t *base, uint8_t *owner, - uint8_t *target, uint64_t target_len, - uint8_t *suffix, uint64_t suffix_len, + uint8_t *patterns, uint64_t patterns_len, bool case_insensitive) { cl_platform_id plat; cl_device_id dev = select_device(id, &plat); @@ -207,8 +206,6 @@ void *gpu_grind_init(int id, uint8_t *base, uint8_t *owner, c->local = clamp_local(dev, GRIND_LOCAL); c->global = (size_t)compute_units(dev) * GRIND_WAVES * c->local; - c->target_len = (uint32_t)target_len; - c->suffix_len = (uint32_t)suffix_len; c->max_iters = GRIND_ITERS_INIT; c->counts_host = (uint32_t *)malloc(c->global * sizeof(uint32_t)); @@ -225,20 +222,6 @@ void *gpu_grind_init(int id, uint8_t *base, uint8_t *owner, for (int j = 0; j < i; ++j) if (alphabet[j] == alphabet[i]) { match_lut[i] = (uint8_t)j; break; } } - uint8_t target_idx[64], suffix_idx[64]; - for (uint64_t i = 0; i < target_len; ++i) { - uint8_t v = 255; - for (int k = 0; k < 58; ++k) - if ((uint8_t)alphabet[k] == target[i]) { v = match_lut[k]; break; } - target_idx[i] = v; - } - for (uint64_t i = 0; i < suffix_len; ++i) { - uint8_t v = 255; - for (int k = 0; k < 58; ++k) - if ((uint8_t)alphabet[k] == suffix[i]) { v = match_lut[k]; break; } - suffix_idx[i] = v; - } - /* SHA-256 block-1 schedule from owner[16..32] (loop invariant). */ uint32_t W1[64]; { @@ -281,11 +264,13 @@ void *gpu_grind_init(int id, uint8_t *base, uint8_t *owner, c->w0 = buf_copy(c->context, sizeof W0, W0); c->sr7 = buf_copy(c->context, sizeof SR7, SR7); c->w1 = buf_copy(c->context, sizeof W1, W1); + uint8_t ptable[VANITY_PT_SIZE]; + vanity_unpack_patterns(patterns, patterns_len, ptable); + c->glyph = buf_copy(c->context, sizeof glyph, glyph); - c->mlut = buf_copy(c->context, sizeof match_lut, match_lut); - c->target = buf_copy(c->context, target_len, target_idx); - c->suffix = buf_copy(c->context, suffix_len, suffix_idx); - c->out = clCreateBuffer(c->context, CL_MEM_WRITE_ONLY, 16, NULL, &err); CK(err, "buf out"); + c->mlut = buf_copy(c->context, sizeof match_lut, match_lut); + c->patterns = buf_copy(c->context, VANITY_PT_SIZE, ptable); + c->out = clCreateBuffer(c->context, CL_MEM_WRITE_ONLY, 16, NULL, &err); CK(err, "buf out"); c->done = clCreateBuffer(c->context, CL_MEM_READ_WRITE, sizeof(cl_int), NULL, &err); CK(err, "buf done"); c->counts = clCreateBuffer(c->context, CL_MEM_WRITE_ONLY, c->global * sizeof(cl_uint), NULL, &err); CK(err, "buf counts"); @@ -295,13 +280,10 @@ void *gpu_grind_init(int id, uint8_t *base, uint8_t *owner, CK(clSetKernelArg(c->kernel, 3, sizeof(cl_mem), &c->w1), "arg w1"); CK(clSetKernelArg(c->kernel, 4, sizeof(cl_mem), &c->glyph), "arg glyph"); CK(clSetKernelArg(c->kernel, 5, sizeof(cl_mem), &c->mlut), "arg mlut"); - CK(clSetKernelArg(c->kernel, 6, sizeof(cl_mem), &c->target), "arg target"); - CK(clSetKernelArg(c->kernel, 7, sizeof(cl_uint), &c->target_len), "arg target_len"); - CK(clSetKernelArg(c->kernel, 8, sizeof(cl_mem), &c->suffix), "arg suffix"); - CK(clSetKernelArg(c->kernel, 9, sizeof(cl_uint), &c->suffix_len), "arg suffix_len"); - CK(clSetKernelArg(c->kernel, 10, sizeof(cl_mem), &c->out), "arg out"); - CK(clSetKernelArg(c->kernel, 11, sizeof(cl_mem), &c->done), "arg done"); - CK(clSetKernelArg(c->kernel, 12, sizeof(cl_mem), &c->counts),"arg counts"); + CK(clSetKernelArg(c->kernel, 6, sizeof(cl_mem), &c->patterns), "arg patterns"); + CK(clSetKernelArg(c->kernel, 7, sizeof(cl_mem), &c->out), "arg out"); + CK(clSetKernelArg(c->kernel, 8, sizeof(cl_mem), &c->done), "arg done"); + CK(clSetKernelArg(c->kernel, 9, sizeof(cl_mem), &c->counts),"arg counts"); return c; } @@ -316,7 +298,7 @@ void gpu_grind_launch(void *opaque, uint8_t *seed) { CK(clEnqueueWriteBuffer(c->queue, c->out, CL_FALSE, 0, 16, out_zero, 0, NULL, NULL), "clear out"); CK(clEnqueueWriteBuffer(c->queue, c->done, CL_FALSE, 0, sizeof zero, &zero, 0, NULL, NULL), "write done"); CK(clSetKernelArg(c->kernel, 0, sizeof(cl_mem), &c->seed), "arg seed"); - CK(clSetKernelArg(c->kernel, 13, sizeof(cl_uint), &c->max_iters), "arg max_iters"); + CK(clSetKernelArg(c->kernel, 10, sizeof(cl_uint), &c->max_iters), "arg max_iters"); c->launch_time = now_sec(); CK(clEnqueueNDRangeKernel(c->queue, c->kernel, 1, NULL, &c->global, &c->local, 0, NULL, &c->event), @@ -356,7 +338,7 @@ void gpu_grind_destroy(void *opaque) { GrindCtx *c = (GrindCtx *)opaque; clFinish(c->queue); if (c->in_flight) clReleaseEvent(c->event); - cl_mem bufs[] = {c->seed,c->w0,c->sr7,c->w1,c->glyph,c->mlut,c->target,c->suffix,c->out,c->done,c->counts}; + cl_mem bufs[] = {c->seed,c->w0,c->sr7,c->w1,c->glyph,c->mlut,c->patterns,c->out,c->done,c->counts}; for (size_t i = 0; i < sizeof bufs / sizeof bufs[0]; ++i) clReleaseMemObject(bufs[i]); clReleaseKernel(c->kernel); clReleaseProgram(c->program); @@ -373,9 +355,8 @@ typedef struct { cl_command_queue queue; cl_program program; cl_kernel kernel; - cl_mem seed, mlut, prefix, suffix, out, done, counts, comb; + cl_mem seed, mlut, patterns, out, done, counts, comb; size_t local, global; - uint32_t prefix_len, suffix_len; uint32_t max_iters; cl_event event; int in_flight; @@ -387,8 +368,8 @@ typedef struct { #define COMB_NIELS_BYTES 160u #define COMB_TABLE_BYTES (52u * 16u * COMB_NIELS_BYTES) -void *gpu_keypair_init(int id, uint8_t *prefix, uint64_t prefix_len, - uint8_t *suffix, uint64_t suffix_len, bool case_insensitive) { +void *gpu_keypair_init(int id, uint8_t *patterns, uint64_t patterns_len, + bool case_insensitive) { cl_platform_id plat; cl_device_id dev = select_device(id, &plat); cl_int err; @@ -406,8 +387,6 @@ void *gpu_keypair_init(int id, uint8_t *prefix, uint64_t prefix_len, c->local = clamp_local(dev, KP_LOCAL); c->global = (size_t)compute_units(dev) * KP_WAVES * c->local; - c->prefix_len = (uint32_t)prefix_len; - c->suffix_len = (uint32_t)suffix_len; c->max_iters = KP_ITERS_INIT; c->counts_host = (uint32_t *)malloc(c->global * sizeof(uint32_t)); @@ -425,24 +404,12 @@ void *gpu_keypair_init(int id, uint8_t *prefix, uint64_t prefix_len, for (int j = 0; j < i; ++j) if (alphabet[j] == alphabet[i]) { match_lut[i] = (uint8_t)j; break; } } - uint8_t prefix_idx[64], suffix_idx[64]; - for (uint64_t i = 0; i < prefix_len; ++i) { - uint8_t v = 255; - for (int k = 0; k < 58; ++k) - if ((uint8_t)alphabet[k] == prefix[i]) { v = match_lut[k]; break; } - prefix_idx[i] = v; - } - for (uint64_t i = 0; i < suffix_len; ++i) { - uint8_t v = 255; - for (int k = 0; k < 58; ++k) - if ((uint8_t)alphabet[k] == suffix[i]) { v = match_lut[k]; break; } - suffix_idx[i] = v; - } + uint8_t ptable[VANITY_PT_SIZE]; + vanity_unpack_patterns(patterns, patterns_len, ptable); - c->seed = clCreateBuffer(c->context, CL_MEM_READ_ONLY, 32, NULL, &err); CK(err, "buf seed"); - c->mlut = buf_copy(c->context, sizeof match_lut, match_lut); - c->prefix = buf_copy(c->context, prefix_len, prefix_idx); - c->suffix = buf_copy(c->context, suffix_len, suffix_idx); + c->seed = clCreateBuffer(c->context, CL_MEM_READ_ONLY, 32, NULL, &err); CK(err, "buf seed"); + c->mlut = buf_copy(c->context, sizeof match_lut, match_lut); + c->patterns = buf_copy(c->context, VANITY_PT_SIZE, ptable); c->out = clCreateBuffer(c->context, CL_MEM_WRITE_ONLY, 32, NULL, &err); CK(err, "buf out"); c->done = clCreateBuffer(c->context, CL_MEM_READ_WRITE, sizeof(cl_int), NULL, &err); CK(err, "buf done"); c->counts = clCreateBuffer(c->context, CL_MEM_WRITE_ONLY, c->global * sizeof(cl_uint), NULL, &err); CK(err, "buf counts"); @@ -461,14 +428,11 @@ void *gpu_keypair_init(int id, uint8_t *prefix, uint64_t prefix_len, } CK(clSetKernelArg(c->kernel, 1, sizeof(cl_mem), &c->mlut), "arg mlut"); - CK(clSetKernelArg(c->kernel, 2, sizeof(cl_mem), &c->prefix), "arg prefix"); - CK(clSetKernelArg(c->kernel, 3, sizeof(cl_uint), &c->prefix_len), "arg prefix_len"); - CK(clSetKernelArg(c->kernel, 4, sizeof(cl_mem), &c->suffix), "arg suffix"); - CK(clSetKernelArg(c->kernel, 5, sizeof(cl_uint), &c->suffix_len), "arg suffix_len"); - CK(clSetKernelArg(c->kernel, 6, sizeof(cl_mem), &c->out), "arg out"); - CK(clSetKernelArg(c->kernel, 7, sizeof(cl_mem), &c->done), "arg done"); - CK(clSetKernelArg(c->kernel, 8, sizeof(cl_mem), &c->counts), "arg counts"); - CK(clSetKernelArg(c->kernel, 9, sizeof(cl_mem), &c->comb), "arg comb"); + CK(clSetKernelArg(c->kernel, 2, sizeof(cl_mem), &c->patterns), "arg patterns"); + CK(clSetKernelArg(c->kernel, 3, sizeof(cl_mem), &c->out), "arg out"); + CK(clSetKernelArg(c->kernel, 4, sizeof(cl_mem), &c->done), "arg done"); + CK(clSetKernelArg(c->kernel, 5, sizeof(cl_mem), &c->counts), "arg counts"); + CK(clSetKernelArg(c->kernel, 6, sizeof(cl_mem), &c->comb), "arg comb"); return c; } @@ -483,7 +447,7 @@ void gpu_keypair_launch(void *opaque, uint8_t *seed) { CK(clEnqueueWriteBuffer(c->queue, c->out, CL_FALSE, 0, 32, out_zero, 0, NULL, NULL), "clear out"); CK(clEnqueueWriteBuffer(c->queue, c->done, CL_FALSE, 0, sizeof zero, &zero, 0, NULL, NULL), "write done"); CK(clSetKernelArg(c->kernel, 0, sizeof(cl_mem), &c->seed), "arg seed"); - CK(clSetKernelArg(c->kernel, 10, sizeof(cl_uint), &c->max_iters), "arg max_iters"); + CK(clSetKernelArg(c->kernel, 7, sizeof(cl_uint), &c->max_iters), "arg max_iters"); c->launch_time = now_sec(); CK(clEnqueueNDRangeKernel(c->queue, c->kernel, 1, NULL, &c->global, &c->local, 0, NULL, &c->event), @@ -523,7 +487,7 @@ void gpu_keypair_destroy(void *opaque) { KeypairCtx *c = (KeypairCtx *)opaque; clFinish(c->queue); if (c->in_flight) clReleaseEvent(c->event); - cl_mem bufs[] = {c->seed,c->mlut,c->prefix,c->suffix,c->out,c->done,c->counts,c->comb}; + cl_mem bufs[] = {c->seed,c->mlut,c->patterns,c->out,c->done,c->counts,c->comb}; for (size_t i = 0; i < sizeof bufs / sizeof bufs[0]; ++i) clReleaseMemObject(bufs[i]); clReleaseKernel(c->kernel); clReleaseProgram(c->program); diff --git a/kernels/opencl/keypair.cl b/kernels/opencl/keypair.cl index fd5e84a..741308e 100644 --- a/kernels/opencl/keypair.cl +++ b/kernels/opencl/keypair.cl @@ -22,8 +22,7 @@ __kernel void vanity_keypair_search( __global const uchar *host_seed, /* 32 bytes */ __constant const uchar *match_lut, /* 58 */ - __global const uchar *prefix, uint prefix_len, - __global const uchar *suffix, uint suffix_len, + __constant const uchar *patterns, __global uchar *out, /* 32 bytes: matched seed */ __global volatile int *done, __global uint *counts, @@ -32,6 +31,18 @@ __kernel void vanity_keypair_search( { ulong idx = get_global_id(0); + uint n_pat = (uint)patterns[0] | ((uint)patterns[1] << 8) + | ((uint)patterns[2] << 16) | ((uint)patterns[3] << 24); + uint prefix_len = 0, suffix_len = 0; + __constant const uchar *prefix = patterns; + __constant const uchar *suffix = patterns; + if (n_pat == 1) { + prefix_len = (uint)patterns[VANITY_PT_PLEN]; + prefix = patterns + VANITY_PT_PREF; + suffix_len = (uint)patterns[VANITY_PT_SLEN]; + suffix = patterns + VANITY_PT_SUF; + } + uchar seed[32]; uchar privatek[64]; uchar pubkey[32]; @@ -108,8 +119,11 @@ __kernel void vanity_keypair_search( | ((uint)pubkey[4*k + 3] ); } - if (fd_base58_check_match_32_words(pubkey_words, prefix, prefix_len, - suffix, suffix_len, match_lut)) { + bool hit = (n_pat <= 1) + ? fd_base58_check_match_32_words(pubkey_words, prefix, prefix_len, + suffix, suffix_len, match_lut) + : fd_base58_check_match_any_32_words(pubkey_words, patterns, match_lut); + if (hit) { if (atomic_cmpxchg(done, 0, 1) == 0) for (int i = 0; i < 32; i++) out[i] = batch_seeds[j][i]; matched = j; diff --git a/kernels/opencl/vanity.cl b/kernels/opencl/vanity.cl index 2d7d385..d7d2b06 100644 --- a/kernels/opencl/vanity.cl +++ b/kernels/opencl/vanity.cl @@ -79,8 +79,7 @@ __kernel void vanity_search( __constant const WORD *W1, /* 64 */ __constant const uchar *glyph, /* 256 */ __constant const uchar *match_lut, /* 58 */ - __global const uchar *target, uint target_len, - __global const uchar *suffix, uint suffix_len, + __constant const uchar *patterns, __global uchar *out, /* 16 bytes: matched seed16 */ __global volatile int *done, __global uint *counts, @@ -88,6 +87,18 @@ __kernel void vanity_search( { ulong idx = get_global_id(0); + uint n_pat = (uint)patterns[0] | ((uint)patterns[1] << 8) + | ((uint)patterns[2] << 16) | ((uint)patterns[3] << 24); + ulong target_len = 0, suffix_len = 0; + __constant const uchar *target = patterns; + __constant const uchar *suffix = patterns; + if (n_pat == 1) { + target_len = (ulong)patterns[VANITY_PT_PLEN]; + target = patterns + VANITY_PT_PREF; + suffix_len = (ulong)patterns[VANITY_PT_SLEN]; + suffix = patterns + VANITY_PT_SUF; + } + /* Bootstrap per-work-item digest: byte-form seed with idx mixed into each 64-bit lane, then packed once into 8 big-endian words. */ WORD digest_words[8]; @@ -122,8 +133,11 @@ __kernel void vanity_search( vanity_pubkey_sha256_words(seed_words, digest_words, W0_fixed, state_r7, W1); - if (fd_base58_check_match_32_words(digest_words, target, target_len, - suffix, suffix_len, match_lut)) { + bool hit = (n_pat <= 1) + ? fd_base58_check_match_32_words(digest_words, target, target_len, + suffix, suffix_len, match_lut) + : fd_base58_check_match_any_32_words(digest_words, patterns, match_lut); + if (hit) { if (atomic_cmpxchg(done, 0, 1) == 0) { for (int k = 0; k < 4; ++k) { WORD w = seed_words[k]; diff --git a/kernels/pattern_table.h b/kernels/pattern_table.h new file mode 100644 index 0000000..775f79a --- /dev/null +++ b/kernels/pattern_table.h @@ -0,0 +1,42 @@ +#ifndef VANITY_PATTERN_TABLE_H +#define VANITY_PATTERN_TABLE_H + +#include +#include + +#define VANITY_MAX_PATTERNS 64 +#define VANITY_MAX_PATTERN_LEN 44 + +/* Flat SoA uploaded to CUDA __constant__ / OpenCL __constant. + Offsets are explicit so host and kernels agree (no struct padding). */ +#define VANITY_PT_N 0 +#define VANITY_PT_PLEN 4 +#define VANITY_PT_SLEN (4 + VANITY_MAX_PATTERNS) +#define VANITY_PT_PREF (VANITY_PT_SLEN + VANITY_MAX_PATTERNS) +#define VANITY_PT_SUF (VANITY_PT_PREF + VANITY_MAX_PATTERNS * VANITY_MAX_PATTERN_LEN) +#define VANITY_PT_SIZE (VANITY_PT_SUF + VANITY_MAX_PATTERNS * VANITY_MAX_PATTERN_LEN) + +/* Packed blob from MatchTargets::gpu_blob → flat SoA. */ +static void vanity_unpack_patterns(const uint8_t *blob, uint64_t blob_len, + uint8_t *out) +{ + memset(out, 0, VANITY_PT_SIZE); + if (blob_len < 6) return; + uint32_t n = 0; + memcpy(&n, blob, 4); + if (n > VANITY_MAX_PATTERNS) n = VANITY_MAX_PATTERNS; + memcpy(out + VANITY_PT_N, &n, 4); + const uint8_t *p = blob + 6; + for (uint32_t i = 0; i < n; i++) { + uint8_t plen = *p++; + out[VANITY_PT_PLEN + i] = plen; + memcpy(out + VANITY_PT_PREF + i * VANITY_MAX_PATTERN_LEN, p, plen); + p += plen; + uint8_t slen = *p++; + out[VANITY_PT_SLEN + i] = slen; + memcpy(out + VANITY_PT_SUF + i * VANITY_MAX_PATTERN_LEN, p, slen); + p += slen; + } +} + +#endif diff --git a/kernels/vanity.cu b/kernels/vanity.cu index 8b6b582..87d58e3 100644 --- a/kernels/vanity.cu +++ b/kernels/vanity.cu @@ -120,14 +120,14 @@ typedef struct { int num_threads; unsigned long long target_cycles; uint64_t out_offset; + int multi; } GpuGrindCtx; extern "C" void* gpu_grind_init( int id, uint8_t *base, uint8_t *owner, - uint8_t *target, uint64_t target_len, - uint8_t *suffix, uint64_t suffix_len, + uint8_t *patterns, uint64_t patterns_len, bool case_insensitive) { cudaSetDevice(id); @@ -159,9 +159,9 @@ extern "C" void* gpu_grind_init( cudaStreamCreate(&ctx->stream); - // Buffer: [seed:32] [base:32] [owner:32] [target_len:8] [target:N] [suffix_len:8] [suffix:M] [out:16] - uint64_t buf_size = 32 + 32 + 32 + 8 + target_len + 8 + suffix_len + 16; - ctx->out_offset = 32 + 32 + 32 + 8 + target_len + 8 + suffix_len; + // Buffer: [seed:32] [base:32] [owner:32] [pat_len:8] [patterns:N] [out:16] + uint64_t buf_size = 32 + 32 + 32 + 8 + patterns_len + 16; + ctx->out_offset = 32 + 32 + 32 + 8 + patterns_len; err = cudaMalloc((void**)&ctx->d_buffer, buf_size); if (err != cudaSuccess) { @@ -191,30 +191,17 @@ extern "C" void* gpu_grind_init( } } - /* Translate ASCII target/suffix to canonical indices in host buffers. */ - uint8_t target_idx[64]; - uint8_t suffix_idx[64]; - for (uint64_t i = 0; i < target_len; ++i) { - uint8_t v = 255; - for (int k = 0; k < 58; ++k) { - if ((uint8_t)alphabet[k] == target[i]) { v = host_match_lut[k]; break; } - } - target_idx[i] = v; - } - for (uint64_t i = 0; i < suffix_len; ++i) { - uint8_t v = 255; - for (int k = 0; k < 58; ++k) { - if ((uint8_t)alphabet[k] == suffix[i]) { v = host_match_lut[k]; break; } - } - suffix_idx[i] = v; - } + uint32_t n_pat = 0; + if (patterns_len >= 4) memcpy(&n_pat, patterns, 4); + ctx->multi = n_pat > 1; + if (ctx->multi) + vanity_upload_pattern_table(patterns, patterns_len); cudaMemcpy(ctx->d_buffer + 32, base, 32, cudaMemcpyHostToDevice); cudaMemcpy(ctx->d_buffer + 64, owner, 32, cudaMemcpyHostToDevice); - cudaMemcpy(ctx->d_buffer + 96, &target_len, 8, cudaMemcpyHostToDevice); - if (target_len > 0) cudaMemcpy(ctx->d_buffer + 104, target_idx, target_len, cudaMemcpyHostToDevice); - cudaMemcpy(ctx->d_buffer + 104 + target_len, &suffix_len, 8, cudaMemcpyHostToDevice); - if (suffix_len > 0) cudaMemcpy(ctx->d_buffer + 104 + target_len + 8, suffix_idx, suffix_len, cudaMemcpyHostToDevice); + cudaMemcpy(ctx->d_buffer + 96, &patterns_len, 8, cudaMemcpyHostToDevice); + if (patterns_len > 0) + cudaMemcpy(ctx->d_buffer + 104, patterns, patterns_len, cudaMemcpyHostToDevice); cudaMemcpyToSymbol(d_match_lut, host_match_lut, sizeof(host_match_lut)); @@ -290,6 +277,9 @@ extern "C" void* gpu_grind_init( return (void *)ctx; } +template +__global__ void vanity_search(uint8_t *buffer, uint64_t stride, unsigned long long max_cycles); + extern "C" void gpu_grind_launch(void *opaque, uint8_t *seed) { GpuGrindCtx *ctx = (GpuGrindCtx *)opaque; @@ -303,10 +293,17 @@ extern "C" void gpu_grind_launch(void *opaque, uint8_t *seed) cudaMemcpyToSymbol(done, &zero, sizeof(int)); cudaMemcpyToSymbol(count, &zero_ull, sizeof(unsigned long long)); - vanity_search<<num_blocks, ctx->num_threads, 0, ctx->stream>>>( - ctx->d_buffer, - (uint64_t)ctx->num_blocks * ctx->num_threads, - ctx->target_cycles); + if (ctx->multi) { + vanity_search<<num_blocks, ctx->num_threads, 0, ctx->stream>>>( + ctx->d_buffer, + (uint64_t)ctx->num_blocks * ctx->num_threads, + ctx->target_cycles); + } else { + vanity_search<<num_blocks, ctx->num_threads, 0, ctx->stream>>>( + ctx->d_buffer, + (uint64_t)ctx->num_blocks * ctx->num_threads, + ctx->target_cycles); + } cudaError_t launch_err = cudaGetLastError(); if (launch_err != cudaSuccess) { @@ -347,18 +344,31 @@ extern "C" void gpu_grind_destroy(void *opaque) // ─── kernel ───────────────────────────────────────────────────────────────── +template __global__ void __launch_bounds__(256, 3) vanity_search(uint8_t *buffer, uint64_t stride, unsigned long long max_cycles) { (void)stride; uint8_t *seed = buffer; - uint64_t target_len; - memcpy(&target_len, buffer + 96, 8); - uint8_t *target = buffer + 104; - uint64_t suffix_len; - memcpy(&suffix_len, buffer + 104 + target_len, 8); - uint8_t *suffix = buffer + 104 + target_len + 8; - uint8_t *out = (buffer + 104 + target_len + suffix_len + 8); + uint64_t pat_len; + memcpy(&pat_len, buffer + 96, 8); + uint8_t *patterns = buffer + 104; + uint8_t *out = buffer + 104 + pat_len; + + /* Hoist the common 1-pattern case so the hot loop calls the original + forceinlined matcher (same register footprint / occupancy). */ + uint32_t n_pat = 0; + memcpy(&n_pat, patterns, 4); + const uint8_t *target = patterns; + ulong target_len = 0; + const uint8_t *suffix = patterns; + ulong suffix_len = 0; + if (n_pat == 1) { + target_len = patterns[6]; + target = patterns + 7; + suffix_len = patterns[7 + target_len]; + suffix = patterns + 8 + target_len; + } uint64_t idx = blockIdx.x * blockDim.x + threadIdx.x; @@ -420,7 +430,10 @@ vanity_search(uint8_t *buffer, uint64_t stride, unsigned long long max_cycles) vanity_pubkey_sha256_words(seed_words, digest_words); - if (fd_base58_check_match_32_words(digest_words, target, target_len, suffix, suffix_len)) + bool hit = MULTI + ? fd_base58_check_match_any_32_words(digest_words) + : fd_base58_check_match_32_words(digest_words, target, target_len, suffix, suffix_len); + if (hit) { if (atomicMax(&done, 1) == 0) { diff --git a/kernels/vanity.h b/kernels/vanity.h index 5779432..02a96d9 100644 --- a/kernels/vanity.h +++ b/kernels/vanity.h @@ -4,12 +4,10 @@ #include #include "utils.h" -extern "C" void* gpu_grind_init(int id, uint8_t *base, uint8_t *owner, uint8_t *target, uint64_t target_len, uint8_t *suffix, uint64_t suffix_len, bool case_insensitive); +extern "C" void* gpu_grind_init(int id, uint8_t *base, uint8_t *owner, uint8_t *patterns, uint64_t patterns_len, bool case_insensitive); extern "C" void gpu_grind_launch(void *ctx, uint8_t *seed); extern "C" int gpu_grind_query(void *ctx); extern "C" void gpu_grind_read(void *ctx, uint8_t *out); extern "C" void gpu_grind_destroy(void *ctx); -__global__ void vanity_search(uint8_t *buffer, uint64_t stride, unsigned long long max_cycles); - #endif diff --git a/kernels/vanity_keypair.cu b/kernels/vanity_keypair.cu index 755de29..7553638 100644 --- a/kernels/vanity_keypair.cu +++ b/kernels/vanity_keypair.cu @@ -1,5 +1,6 @@ #include #include +#include #include #include "vanity_keypair.h" #include "base58.h" @@ -17,8 +18,8 @@ __device__ static unsigned long long kp_count = 0; #define KP_BATCH KP_BATCH_MAX #endif -static __global__ void __launch_bounds__(KP_MAX_THREADS) -vanity_keypair_search(uint8_t *buffer, uint64_t stride, unsigned long long max_cycles, +template +static __global__ void vanity_keypair_search(uint8_t *buffer, uint64_t stride, unsigned long long max_cycles, const ge_niels *comb); /* Device-wide nanosecond clock. clock64() is per-SM and cannot be compared @@ -80,12 +81,12 @@ typedef struct { int num_threads; unsigned long long target_cycles; uint64_t out_offset; + int multi; } GpuKeypairCtx; extern "C" void* gpu_keypair_init( int id, - uint8_t *prefix, uint64_t prefix_len, - uint8_t *suffix, uint64_t suffix_len, + uint8_t *patterns, uint64_t patterns_len, bool case_insensitive) { cudaSetDevice(id); @@ -97,22 +98,32 @@ extern "C" void* gpu_keypair_init( exit(EXIT_FAILURE); } + uint32_t n_pat = 0; + if (patterns_len >= 4) memcpy(&n_pat, patterns, 4); + int multi = n_pat > 1; + if (multi) + vanity_upload_pattern_table(patterns, patterns_len); + int nthreads = KP_MAX_THREADS; - int nblocks = one_wave_blocks(vanity_keypair_search, nthreads, - prop.multiProcessorCount, "gpu_keypair_init"); + int nblocks = multi + ? one_wave_blocks(vanity_keypair_search, nthreads, + prop.multiProcessorCount, "gpu_keypair_init") + : one_wave_blocks(vanity_keypair_search, nthreads, + prop.multiProcessorCount, "gpu_keypair_init"); GpuKeypairCtx *ctx = (GpuKeypairCtx *)malloc(sizeof(GpuKeypairCtx)); ctx->device_id = id; ctx->num_blocks = nblocks; ctx->num_threads = nthreads; + ctx->multi = multi; /* 2s wall-clock slice, timed with %globaltimer (nanoseconds). */ ctx->target_cycles = 2000000000ULL; cudaStreamCreate(&ctx->stream); - // Buffer: [seed:32] [prefix_len:8] [prefix:N] [suffix_len:8] [suffix:M] [out:32] - uint64_t buf_size = 32 + 8 + prefix_len + 8 + suffix_len + 32; - ctx->out_offset = 32 + 8 + prefix_len + 8 + suffix_len; + // Buffer: [seed:32] [pat_len:8] [patterns:N] [out:32] + uint64_t buf_size = 32 + 8 + patterns_len + 32; + ctx->out_offset = 32 + 8 + patterns_len; err = cudaMalloc((void**)&ctx->d_buffer, buf_size); if (err != cudaSuccess) { @@ -139,29 +150,9 @@ extern "C" void* gpu_keypair_init( } } - uint8_t prefix_idx[64]; - uint8_t suffix_idx[64]; - for (uint64_t i = 0; i < prefix_len; ++i) { - uint8_t v = 255; - for (int k = 0; k < 58; ++k) { - if ((uint8_t)alphabet[k] == prefix[i]) { v = host_match_lut[k]; break; } - } - prefix_idx[i] = v; - } - for (uint64_t i = 0; i < suffix_len; ++i) { - uint8_t v = 255; - for (int k = 0; k < 58; ++k) { - if ((uint8_t)alphabet[k] == suffix[i]) { v = host_match_lut[k]; break; } - } - suffix_idx[i] = v; - } - - uint64_t off = 32; - cudaMemcpy(ctx->d_buffer + off, &prefix_len, 8, cudaMemcpyHostToDevice); off += 8; - if (prefix_len > 0) cudaMemcpy(ctx->d_buffer + off, prefix_idx, prefix_len, cudaMemcpyHostToDevice); - off += prefix_len; - cudaMemcpy(ctx->d_buffer + off, &suffix_len, 8, cudaMemcpyHostToDevice); off += 8; - if (suffix_len > 0) cudaMemcpy(ctx->d_buffer + off, suffix_idx, suffix_len, cudaMemcpyHostToDevice); + cudaMemcpy(ctx->d_buffer + 32, &patterns_len, 8, cudaMemcpyHostToDevice); + if (patterns_len > 0) + cudaMemcpy(ctx->d_buffer + 40, patterns, patterns_len, cudaMemcpyHostToDevice); cudaMemcpyToSymbol(d_match_lut, host_match_lut, sizeof(host_match_lut)); } @@ -186,11 +177,19 @@ extern "C" void gpu_keypair_launch(void *opaque, uint8_t *seed) cudaMemcpyToSymbol(kp_done, &zero, sizeof(int)); cudaMemcpyToSymbol(kp_count, &zero_ull, sizeof(unsigned long long)); - vanity_keypair_search<<num_blocks, ctx->num_threads, 0, ctx->stream>>>( - ctx->d_buffer, - (uint64_t)ctx->num_blocks * ctx->num_threads, - ctx->target_cycles, - ctx->d_comb); + if (ctx->multi) { + vanity_keypair_search<<num_blocks, ctx->num_threads, 0, ctx->stream>>>( + ctx->d_buffer, + (uint64_t)ctx->num_blocks * ctx->num_threads, + ctx->target_cycles, + ctx->d_comb); + } else { + vanity_keypair_search<<num_blocks, ctx->num_threads, 0, ctx->stream>>>( + ctx->d_buffer, + (uint64_t)ctx->num_blocks * ctx->num_threads, + ctx->target_cycles, + ctx->d_comb); + } cudaError_t launch_err = cudaGetLastError(); if (launch_err != cudaSuccess) { @@ -238,6 +237,7 @@ static __device__ __forceinline__ void kp_sha512_32(const unsigned char seed[32] sha512_32(seed, out); } +template static __global__ void __launch_bounds__(KP_MAX_THREADS) vanity_keypair_search(uint8_t *buffer, uint64_t stride, unsigned long long max_cycles, const ge_niels *comb) @@ -245,15 +245,23 @@ vanity_keypair_search(uint8_t *buffer, uint64_t stride, unsigned long long max_c (void)stride; uint8_t *host_seed = buffer; - uint64_t prefix_len; - memcpy(&prefix_len, buffer + 32, 8); - uint8_t *prefix = buffer + 40; - - uint64_t suffix_len; - memcpy(&suffix_len, buffer + 40 + prefix_len, 8); - uint8_t *suffix = buffer + 40 + prefix_len + 8; - - uint8_t *out = suffix + suffix_len; + uint64_t pat_len; + memcpy(&pat_len, buffer + 32, 8); + uint8_t *patterns = buffer + 40; + uint8_t *out = buffer + 40 + pat_len; + + uint32_t n_pat = 0; + memcpy(&n_pat, patterns, 4); + const uint8_t *prefix = patterns; + ulong prefix_len = 0; + const uint8_t *suffix = patterns; + ulong suffix_len = 0; + if (n_pat == 1) { + prefix_len = patterns[6]; + prefix = patterns + 7; + suffix_len = patterns[7 + prefix_len]; + suffix = patterns + 8 + prefix_len; + } uint64_t idx = blockIdx.x * blockDim.x + threadIdx.x; @@ -325,7 +333,10 @@ vanity_keypair_search(uint8_t *buffer, uint64_t stride, unsigned long long max_c | ((uint)pubkey[4*k + 3] ); } - if (fd_base58_check_match_32_words(pubkey_words, prefix, prefix_len, suffix, suffix_len)) + bool hit = MULTI + ? fd_base58_check_match_any_32_words(pubkey_words) + : fd_base58_check_match_32_words(pubkey_words, prefix, prefix_len, suffix, suffix_len); + if (hit) { if (atomicMax(&kp_done, 1) == 0) { memcpy(out, batch_seeds[j], 32); diff --git a/kernels/vanity_keypair.h b/kernels/vanity_keypair.h index a83d4d8..af0ead8 100644 --- a/kernels/vanity_keypair.h +++ b/kernels/vanity_keypair.h @@ -3,7 +3,7 @@ #include -extern "C" void* gpu_keypair_init(int id, uint8_t *prefix, uint64_t prefix_len, uint8_t *suffix, uint64_t suffix_len, bool case_insensitive); +extern "C" void* gpu_keypair_init(int id, uint8_t *patterns, uint64_t patterns_len, bool case_insensitive); extern "C" void gpu_keypair_launch(void *ctx, uint8_t *seed); extern "C" int gpu_keypair_query(void *ctx); extern "C" void gpu_keypair_read(void *ctx, uint8_t *out); diff --git a/src/fast/check_match.rs b/src/fast/check_match.rs index cdb4e21..faeb27c 100644 --- a/src/fast/check_match.rs +++ b/src/fast/check_match.rs @@ -3,6 +3,7 @@ use fd_bs58::constants::{ }; pub const MAX_PATTERN_LEN: usize = 44; +pub const MAX_PATTERNS: usize = 64; const ALPHABET: &[u8; 58] = b"123456789ABCDEFGHJKLMNPQRSTUVWXYZabcdefghijkmnopqrstuvwxyz"; @@ -49,20 +50,36 @@ pub fn match_lut(case_insensitive: bool) -> &'static [u8; 58] { } } -pub struct MatchTarget { +struct PatternIdx { prefix_idx: [u8; MAX_PATTERN_LEN], prefix_len: u8, suffix_idx: [u8; MAX_PATTERN_LEN], suffix_len: u8, +} + +/// One or more prefix/suffix pairs. A pubkey matches if it matches any pair. +/// +/// GPU blob layout (little-endian), shared with CUDA / OpenCL: +/// ```text +/// u32 n +/// u8 max_prefix_len +/// u8 max_suffix_len +/// n times: u8 plen, plen bytes, u8 slen, slen bytes +/// ``` +/// Bytes are canonical raw_base58 indices, not ASCII. +pub struct MatchTargets { + patterns: Vec, + max_prefix_len: u8, + max_suffix_len: u8, match_lut: &'static [u8; 58], } -impl MatchTarget { - pub fn new( - prefix: &str, - suffix: &str, - case_insensitive: bool, - ) -> Self { +impl MatchTargets { + pub fn new(patterns: &[(&str, &str)], case_insensitive: bool) -> Self { + assert!( + patterns.len() <= MAX_PATTERNS, + "at most {MAX_PATTERNS} patterns" + ); let alphabet = if case_insensitive { ALPHABET_CI } else { @@ -70,37 +87,64 @@ impl MatchTarget { }; let lut = match_lut(case_insensitive); - let mut prefix_idx = [0u8; MAX_PATTERN_LEN]; - debug_assert!(prefix.len() <= MAX_PATTERN_LEN); - for (i, &b) in prefix.as_bytes().iter().enumerate() { - prefix_idx[i] = char_to_canonical(b, alphabet, lut); - } - - let mut suffix_idx = [0u8; MAX_PATTERN_LEN]; - debug_assert!(suffix.len() <= MAX_PATTERN_LEN); - for (i, &b) in suffix.as_bytes().iter().enumerate() { - suffix_idx[i] = char_to_canonical(b, alphabet, lut); + let mut out = Vec::with_capacity(patterns.len()); + let mut max_prefix_len = 0u8; + let mut max_suffix_len = 0u8; + for &(prefix, suffix) in patterns { + debug_assert!(prefix.len() <= MAX_PATTERN_LEN); + debug_assert!(suffix.len() <= MAX_PATTERN_LEN); + let mut prefix_idx = [0u8; MAX_PATTERN_LEN]; + for (i, &b) in prefix.as_bytes().iter().enumerate() { + prefix_idx[i] = char_to_canonical(b, alphabet, lut); + } + let mut suffix_idx = [0u8; MAX_PATTERN_LEN]; + for (i, &b) in suffix.as_bytes().iter().enumerate() { + suffix_idx[i] = char_to_canonical(b, alphabet, lut); + } + let plen = prefix.len() as u8; + let slen = suffix.len() as u8; + max_prefix_len = max_prefix_len.max(plen); + max_suffix_len = max_suffix_len.max(slen); + out.push(PatternIdx { + prefix_idx, + prefix_len: plen, + suffix_idx, + suffix_len: slen, + }); } Self { - prefix_idx, - prefix_len: prefix.len() as u8, - suffix_idx, - suffix_len: suffix.len() as u8, + patterns: out, + max_prefix_len, + max_suffix_len, match_lut: lut, } } #[inline] pub fn matches(&self, bytes: &[u8; 32]) -> bool { - check_match_32( - bytes, - &self.prefix_idx, - self.prefix_len, - &self.suffix_idx, - self.suffix_len, - self.match_lut, - ) + check_match_32(bytes, &self.patterns, self.match_lut) + } + + /// Canonical-index blob consumed by `gpu_grind_init` / `gpu_keypair_init`. + pub fn gpu_blob(&self) -> Vec { + let mut v = Vec::with_capacity( + 6 + self + .patterns + .iter() + .map(|p| 2 + p.prefix_len as usize + p.suffix_len as usize) + .sum::(), + ); + v.extend_from_slice(&(self.patterns.len() as u32).to_le_bytes()); + v.push(self.max_prefix_len); + v.push(self.max_suffix_len); + for p in &self.patterns { + v.push(p.prefix_len); + v.extend_from_slice(&p.prefix_idx[..p.prefix_len as usize]); + v.push(p.suffix_len); + v.extend_from_slice(&p.suffix_idx[..p.suffix_len as usize]); + } + v } } @@ -148,10 +192,7 @@ fn ensure_limb( #[inline] fn check_match_32( bytes: &[u8; 32], - prefix_idx: &[u8; MAX_PATTERN_LEN], - prefix_len: u8, - suffix_idx: &[u8; MAX_PATTERN_LEN], - suffix_len: u8, + patterns: &[PatternIdx], match_lut: &[u8; 58], ) -> bool { let mut in_leading_0s = 0usize; @@ -205,40 +246,52 @@ fn check_match_32( let skip = raw_leading_0s - in_leading_0s; let encoded_length = RAW58_SZ_32 - skip; - for i in 0..prefix_len as usize { - let target = prefix_idx[i]; - let rb_idx = skip + i; - ensure_limb( - &intermediate, - &mut raw_base58, - &mut limbs_done, - rb_idx / 5, - ); - if match_lut[raw_base58[rb_idx] as usize] != target { - return false; - } + if patterns.is_empty() { + return true; } - if suffix_len > 0 { - let suffix_len = suffix_len as usize; - let tail_start = skip + encoded_length - suffix_len; - let last_limb = (skip + encoded_length - 1) / 5; - ensure_limb( - &intermediate, - &mut raw_base58, - &mut limbs_done, - last_limb, - ); - for i in 0..suffix_len { - let target = suffix_idx[i]; - if match_lut[raw_base58[tail_start + i] as usize] != target + // Same per-char early-reject walk as the original single-pattern + // matcher; extra patterns re-use already-emitted limbs. + for p in patterns { + let mut ok = true; + for i in 0..p.prefix_len as usize { + let rb_idx = skip + i; + ensure_limb( + &intermediate, + &mut raw_base58, + &mut limbs_done, + rb_idx / 5, + ); + if match_lut[raw_base58[rb_idx] as usize] != p.prefix_idx[i] { - return false; + ok = false; + break; } } + if ok && p.suffix_len > 0 { + let suffix_len = p.suffix_len as usize; + let tail_start = skip + encoded_length - suffix_len; + let last_limb = (skip + encoded_length - 1) / 5; + ensure_limb( + &intermediate, + &mut raw_base58, + &mut limbs_done, + last_limb, + ); + for i in 0..suffix_len { + if match_lut[raw_base58[tail_start + i] as usize] + != p.suffix_idx[i] + { + ok = false; + break; + } + } + } + if ok { + return true; + } } - - true + false } #[cfg(test)] @@ -255,16 +308,16 @@ mod tests { ) -> bool { let s = fd_bs58::encode_32(*bytes); if ci { - let lc: String = - s.chars() - .map(|c| { - if c == 'L' { - c - } else { - c.to_ascii_lowercase() - } - }) - .collect(); + let lc: String = s + .chars() + .map(|c| { + if c == 'L' { + c + } else { + c.to_ascii_lowercase() + } + }) + .collect(); (prefix.is_empty() || lc.starts_with(prefix)) && (suffix.is_empty() || lc.ends_with(suffix)) } else { @@ -275,7 +328,7 @@ mod tests { #[test] fn agrees_with_full_encode_random() { - let target = MatchTarget::new("zz", "LM", false); + let target = MatchTargets::new(&[("zz", "LM")], false); for i in 0u32..5000 { let h: [u8; 64] = Sha512::digest(i.to_le_bytes()).into(); let seed: [u8; 32] = h[..32].try_into().unwrap(); @@ -290,7 +343,7 @@ mod tests { #[test] fn agrees_with_full_encode_ci() { - let target = MatchTarget::new("mithriL", "", true); + let target = MatchTargets::new(&[("mithriL", "")], true); for i in 0u32..2000 { let h: [u8; 64] = Sha512::digest((i ^ 0xdeadbeef).to_le_bytes()).into(); @@ -315,8 +368,32 @@ mod tests { let bytes = fd_bs58::decode_32(key).unwrap(); let prefix = &key[..key.len().min(3)]; let suffix = &key[key.len().saturating_sub(2)..]; - let target = MatchTarget::new(prefix, suffix, false); + let target = MatchTargets::new(&[(prefix, suffix)], false); assert!(target.matches(&bytes), "{key}"); } } + + #[test] + fn or_of_two_patterns() { + let key = "XkCriyrNwS3G4rzAXtG5B1nnvb5Ka1JtCku93VqeKAr"; + let bytes = fd_bs58::decode_32(key).unwrap(); + let miss = MatchTargets::new(&[("zzz", ""), ("aaa", "qq")], false); + assert!(!miss.matches(&bytes)); + let hit = + MatchTargets::new(&[("zzz", ""), ("XkC", "KAr")], false); + assert!(hit.matches(&bytes)); + let pref_only = + MatchTargets::new(&[("nope", "nope"), ("XkC", "")], false); + assert!(pref_only.matches(&bytes)); + } + + #[test] + fn gpu_blob_roundtrip_header() { + let t = MatchTargets::new(&[("ab", "xy"), ("Z", "")], false); + let blob = t.gpu_blob(); + assert_eq!(&blob[..4], &2u32.to_le_bytes()); + assert_eq!(blob[4], 2); // max prefix + assert_eq!(blob[5], 2); // max suffix + assert_eq!(blob[6], 2); // plen of first + } } diff --git a/src/fast/mod.rs b/src/fast/mod.rs index b6f3c32..aa72f02 100644 --- a/src/fast/mod.rs +++ b/src/fast/mod.rs @@ -4,7 +4,7 @@ mod group; pub mod sha512_simd; pub mod simd; -use check_match::MatchTarget; +pub use check_match::{MatchTargets, MAX_PATTERNS, MAX_PATTERN_LEN}; use field::{batch_invert, Fe}; use group::{edwards_d2, Niels, Point}; @@ -195,7 +195,7 @@ unsafe fn keygen_batch_simd( #[cfg(target_arch = "x86_64")] #[target_feature(enable = "avx512f,avx512ifma,avx512dq")] -unsafe fn grind_thread_simd(target: &MatchTarget, count: u32) { +unsafe fn grind_thread_simd(target: &MatchTargets, count: u32) { check_write_permissions(); let mut seeds: [[u8; 32]; BATCH] = [[0u8; 32]; BATCH]; for s in seeds.iter_mut() { @@ -230,7 +230,7 @@ unsafe fn grind_thread_simd(target: &MatchTarget, count: u32) { TOTAL_ATTEMPTS.fetch_add(local, Ordering::Relaxed); } -fn grind_thread_scalar(target: &MatchTarget, count: u32) { +fn grind_thread_scalar(target: &MatchTargets, count: u32) { check_write_permissions(); let mut seeds: [[u8; 32]; BATCH] = [[0u8; 32]; BATCH]; for s in seeds.iter_mut() { @@ -321,15 +321,7 @@ pub fn backend_name() -> &'static str { /// Run batched CPU keypair workers until `count` matches or abort. /// Call [`reset_grind`] first if coordinating with a GPU thread that shares /// these counters via [`add_attempts`] / [`note_found`] / [`is_done`]. -pub fn run_cpu_workers( - prefix: &'static str, - suffix: &'static str, - case_insensitive: bool, - num_cpus: u32, - count: u32, -) { - let target = MatchTarget::new(prefix, suffix, case_insensitive); - +pub fn run_cpu_workers(target: &MatchTargets, num_cpus: u32, count: u32) { (0..num_cpus) .into_par_iter() .for_each(|_| { diff --git a/src/main.rs b/src/main.rs index 46a41d5..5cca70a 100644 --- a/src/main.rs +++ b/src/main.rs @@ -47,6 +47,24 @@ pub enum Command { Deploy(DeployArgs), } +/// Repeatable vanity target: a pubkey matches if it matches any `--pattern`. +/// +/// Syntax (`.` is not in base58, so `...` is an unambiguous split): +/// `--pattern Cavey...CooL` prefix `Cavey` and suffix `CooL` +/// `--pattern Harmonic...` prefix only +/// `--pattern ...pump` suffix only +/// `--pattern Harmonic` prefix only (the `...` may be omitted) +#[derive(Debug, Parser)] +pub struct PatternArgs { + /// Vanity pattern. Repeatable; OR'd together. + #[clap(long, action = clap::ArgAction::Append)] + pub pattern: Vec, + + /// Whether user cares about the case of the pubkey + #[clap(long, default_value_t = false)] + pub case_insensitive: bool, +} + #[derive(Debug, Parser)] pub struct GrindArgs { /// The pubkey that will be the signer for the CreateAccountWithSeed instruction @@ -57,16 +75,8 @@ pub struct GrindArgs { #[clap(long, value_parser = parse_pubkey)] pub owner: Pubkey, - /// The target prefix for the pubkey - #[clap(long)] - pub prefix: Option, - - #[clap(long)] - pub suffix: Option, - - /// Whether user cares about the case of the pubkey - #[clap(long, default_value_t = false)] - pub case_insensitive: bool, + #[clap(flatten)] + pub spec: PatternArgs, /// Number of gpus to use for mining #[clap(long, default_value_t = 1)] @@ -84,17 +94,8 @@ pub struct GrindArgs { #[derive(Debug, Parser)] pub struct GrindKeypairArgs { - /// The target prefix for the pubkey - #[clap(long)] - pub prefix: Option, - - /// The target suffix for the pubkey - #[clap(long)] - pub suffix: Option, - - /// Whether user cares about the case of the pubkey - #[clap(long, default_value_t = false)] - pub case_insensitive: bool, + #[clap(flatten)] + pub spec: PatternArgs, /// Number of gpus to use for mining #[clap(long, default_value_t = 1)] @@ -323,6 +324,7 @@ fn bs58_probability( } } +#[allow(dead_code)] fn expected_attempts( prefix: &str, suffix: &str, @@ -336,6 +338,147 @@ fn expected_attempts( } } +// ─── --pattern ────────────────────────────────────────────────────────────── + +#[derive(Clone, Debug, PartialEq, Eq)] +struct VanityPattern { + prefix: String, + suffix: String, +} + +#[derive(Clone, Debug)] +struct VanitySpec { + patterns: Vec, + case_insensitive: bool, +} + +impl VanitySpec { + fn pairs(&self) -> Vec<(&str, &str)> { + self.patterns + .iter() + .map(|p| (p.prefix.as_str(), p.suffix.as_str())) + .collect() + } + + fn matches(&self, pubkey: &str) -> bool { + self.patterns.iter().any(|p| { + matches_target( + pubkey, + &p.prefix, + &p.suffix, + self.case_insensitive, + ) + }) + } + + fn label(&self) -> String { + self.patterns + .iter() + .map(|p| format_target_label(&p.prefix, &p.suffix)) + .collect::>() + .join(" | ") + } + + fn probability(&self) -> f64 { + self.patterns + .iter() + .map(|p| { + bs58_probability( + &p.prefix, + &p.suffix, + self.case_insensitive, + ) + }) + .sum::() + .min(1.0) + } + + fn expected_attempts(&self) -> f64 { + let p = self.probability(); + if p <= 0.0 { + f64::INFINITY + } else { + 1.0 / p + } + } +} + +/// Split `Prefix...Suffix`, `Prefix...`, `...Suffix`, or a bare prefix. +fn parse_pattern(raw: &str) -> VanityPattern { + let (prefix, suffix) = match raw.split_once("...") { + Some((p, s)) => { + assert!( + !s.contains("..."), + "--pattern may contain only one '...' separator" + ); + (p.to_string(), s.to_string()) + } + None => (raw.to_string(), String::new()), + }; + assert!( + !prefix.is_empty() || !suffix.is_empty(), + "--pattern must include a prefix and/or suffix (e.g. Harmonic... or ...CooL or Cavey...CooL)" + ); + assert!( + prefix.len() <= fast::MAX_PATTERN_LEN + && suffix.len() <= fast::MAX_PATTERN_LEN, + "--pattern prefix/suffix each at most {} characters", + fast::MAX_PATTERN_LEN + ); + VanityPattern { prefix, suffix } +} + +/// `q` matches every string that `p` matches (and possibly more). +fn dominates(q: &VanityPattern, p: &VanityPattern) -> bool { + p.prefix.starts_with(&q.prefix) && p.suffix.ends_with(&q.suffix) +} + +fn prune_dominated(patterns: &[VanityPattern]) -> Vec { + let mut out = Vec::new(); + for p in patterns { + if out.iter().any(|q| dominates(q, p)) { + continue; + } + out.retain(|q| !dominates(p, q)); + out.push(p.clone()); + } + out +} + +fn resolve_spec(args: &PatternArgs) -> VanitySpec { + let raw = &args.pattern; + assert!( + !raw.is_empty(), + "supply at least one --pattern" + ); + assert!( + raw.len() <= fast::MAX_PATTERNS, + "at most {} --pattern flags", + fast::MAX_PATTERNS + ); + + let mut patterns = Vec::with_capacity(raw.len()); + for r in &raw { + let pat = parse_pattern(r); + validate_bs58("pattern prefix", &pat.prefix); + validate_bs58("pattern suffix", &pat.suffix); + patterns.push(VanityPattern { + prefix: maybe_bs58_aware_lowercase( + &pat.prefix, + args.case_insensitive, + ), + suffix: maybe_bs58_aware_lowercase( + &pat.suffix, + args.case_insensitive, + ), + }); + } + VanitySpec { + patterns: prune_dominated(&patterns), + case_insensitive: args.case_insensitive, + } +} + // ─── formatting ───────────────────────────────────────────────────────────── fn format_duration(secs: f64) -> String { @@ -635,41 +778,35 @@ pub fn deploy_with_max_program_len_with_seed( fn grind(mut args: GrindArgs) { maybe_update_num_cpus(&mut args.num_cpus); - let prefix = get_validated_bs58( - "prefix", - &args.prefix, - args.case_insensitive, - ); - let suffix = get_validated_bs58( - "suffix", - &args.suffix, - args.case_insensitive, - ); - - let expected = - expected_attempts(prefix, suffix, args.case_insensitive); - let prob = bs58_probability(prefix, suffix, args.case_insensitive); + let spec = resolve_spec(&args.spec); + let expected = spec.expected_attempts(); + let prob = spec.probability(); #[cfg(feature = "gpu")] eprintln!("using {} cpus, {} gpus", args.num_cpus, args.num_gpus); #[cfg(not(feature = "gpu"))] eprintln!("using {} cpus", args.num_cpus); - let target_label = format_target_label(prefix, suffix); eprintln!( "target: {} | probability: {:.6e} | expected: {} attempts", - target_label, + spec.label(), prob, (expected as u64).to_formatted_string(&Locale::en) ); let target_count = args.count; let shutdown = Arc::new(AtomicBool::new(false)); + let spec = Arc::new(spec); #[cfg(feature = "gpu")] let gpu_thread = if args.num_gpus > 0 { let num_gpus = args.num_gpus; let base = args.base; let owner = args.owner; - let ci = args.case_insensitive; + let pairs = spec.pairs(); + let blob = Arc::new( + fast::MatchTargets::new(&pairs, spec.case_insensitive) + .gpu_blob(), + ); + let spec = Arc::clone(&spec); Some( thread::Builder::new() .name("gpu_mgr".into()) @@ -687,11 +824,9 @@ fn grind(mut args: GrindArgs) { id as i32, base.as_ref().as_ptr(), owner.as_ref().as_ptr(), - prefix.as_ptr(), - prefix.len() as u64, - suffix.as_ptr(), - suffix.len() as u64, - ci, + blob.as_ptr(), + blob.len() as u64, + spec.case_insensitive, ) }, || done(target_count), @@ -703,12 +838,8 @@ fn grind(mut args: GrindArgs) { .finalize() .into(); let out_str = fd_bs58::encode_32(reconstructed); - let out_str_check = - maybe_bs58_aware_lowercase(&out_str, ci); - if out_str_check.starts_with(prefix) - && out_str_check.ends_with(suffix) - { + if spec.matches(&out_str) { eprintln!( "\r\x1b[Kgpu {} match: {} in {:.3}s", i, out_str, time_sec @@ -769,7 +900,7 @@ fn grind(mut args: GrindArgs) { local_batch -= 4096; } - if matches_target(&pubkey, prefix, suffix, args.case_insensitive) { + if spec.matches(&pubkey) { if local_batch > 0 { TOTAL_ATTEMPTS.fetch_add(local_batch, Ordering::Relaxed); local_batch = 0; @@ -819,20 +950,9 @@ fn grind(mut args: GrindArgs) { fn grind_keypair(mut args: GrindKeypairArgs) { check_write_permissions(); maybe_update_num_cpus(&mut args.num_cpus); - let prefix = get_validated_bs58( - "prefix", - &args.prefix, - args.case_insensitive, - ); - let suffix = get_validated_bs58( - "suffix", - &args.suffix, - args.case_insensitive, - ); - - let expected = - expected_attempts(prefix, suffix, args.case_insensitive); - let prob = bs58_probability(prefix, suffix, args.case_insensitive); + let spec = resolve_spec(&args.spec); + let expected = spec.expected_attempts(); + let prob = spec.probability(); #[cfg(feature = "gpu")] eprintln!( "using {} cpus, {} gpus (cpu backend: {})", @@ -846,10 +966,9 @@ fn grind_keypair(mut args: GrindKeypairArgs) { args.num_cpus, fast::backend_name() ); - let target_label = format_target_label(prefix, suffix); eprintln!( "target: {} | probability: {:.6e} | expected: {} attempts", - target_label, + spec.label(), prob, (expected as u64).to_formatted_string(&Locale::en) ); @@ -858,6 +977,12 @@ fn grind_keypair(mut args: GrindKeypairArgs) { fast::reset_grind(); let shutdown = Arc::new(AtomicBool::new(false)); let grind_start = Instant::now(); + let match_targets = { + let pairs = spec.pairs(); + fast::MatchTargets::new(&pairs, spec.case_insensitive) + }; + #[cfg(feature = "gpu")] + let spec = Arc::new(spec); // Reporter reads the shared fast-path counters (CPU + GPU both update them). let reporter = { @@ -880,7 +1005,8 @@ fn grind_keypair(mut args: GrindKeypairArgs) { #[cfg(feature = "gpu")] let gpu_thread = if args.num_gpus > 0 { let num_gpus = args.num_gpus; - let ci = args.case_insensitive; + let blob = Arc::new(match_targets.gpu_blob()); + let spec = Arc::clone(&spec); Some( thread::Builder::new() .name("gpu_mgr".into()) @@ -896,11 +1022,9 @@ fn grind_keypair(mut args: GrindKeypairArgs) { |id| unsafe { gpu_keypair_init( id as i32, - prefix.as_ptr(), - prefix.len() as u64, - suffix.as_ptr(), - suffix.len() as u64, - ci, + blob.as_ptr(), + blob.len() as u64, + spec.case_insensitive, ) }, || { @@ -915,12 +1039,8 @@ fn grind_keypair(mut args: GrindKeypairArgs) { let pubkey_bytes = signing_key.verifying_key().to_bytes(); let pubkey_str = fd_bs58::encode_32(pubkey_bytes); - let pubkey_check = - maybe_bs58_aware_lowercase(&pubkey_str, ci); - if pubkey_check.starts_with(prefix) - && pubkey_check.ends_with(suffix) - { + if spec.matches(&pubkey_str) { let prev = fast::note_found(); if prev < target_count { eprintln!( @@ -944,13 +1064,7 @@ fn grind_keypair(mut args: GrindKeypairArgs) { None }; - fast::run_cpu_workers( - prefix, - suffix, - args.case_insensitive, - args.num_cpus, - target_count, - ); + fast::run_cpu_workers(&match_targets, args.num_cpus, target_count); // CPU workers finished (found enough or aborted); stop GPU too. fast::request_abort(); @@ -1277,24 +1391,15 @@ fn check_write_permissions() { } } -fn get_validated_bs58( - label: &str, - value: &Option, - case_insensitive: bool, -) -> &'static str { +fn validate_bs58(label: &str, value: &str) { const BS58_CHARS: &str = "123456789ABCDEFGHJKLMNPQRSTUVWXYZabcdefghijkmnopqrstuvwxyz"; - if let Some(ref s) = value { - for c in s.chars() { - assert!( - BS58_CHARS.contains(c), - "your {label} contains invalid bs58: {c}" - ); - } - let validated = maybe_bs58_aware_lowercase(s, case_insensitive); - return validated.leak(); + for c in value.chars() { + assert!( + BS58_CHARS.contains(c), + "your {label} contains invalid bs58: {c}" + ); } - "" } fn maybe_bs58_aware_lowercase( @@ -1359,10 +1464,8 @@ extern "C" { id: i32, base: *const u8, owner: *const u8, - target: *const u8, - target_len: u64, - suffix: *const u8, - suffix_len: u64, + patterns: *const u8, + patterns_len: u64, case_insensitive: bool, ) -> *mut std::ffi::c_void; pub fn gpu_grind_launch( @@ -1375,10 +1478,8 @@ extern "C" { pub fn gpu_keypair_init( id: i32, - prefix: *const u8, - prefix_len: u64, - suffix: *const u8, - suffix_len: u64, + patterns: *const u8, + patterns_len: u64, case_insensitive: bool, ) -> *mut std::ffi::c_void; pub fn gpu_keypair_launch( @@ -1442,4 +1543,67 @@ mod tests { fn bs58_ci_factor_skips_non_letters() { assert_eq!(bs58_ci_factor("1A", ""), 2.0); } + + #[test] + fn parse_pattern_splits_on_ellipsis() { + assert_eq!( + parse_pattern("Cavey...CooL"), + VanityPattern { + prefix: "Cavey".into(), + suffix: "CooL".into(), + } + ); + assert_eq!( + parse_pattern("Harmonic..."), + VanityPattern { + prefix: "Harmonic".into(), + suffix: String::new(), + } + ); + assert_eq!( + parse_pattern("...pump"), + VanityPattern { + prefix: String::new(), + suffix: "pump".into(), + } + ); + assert_eq!( + parse_pattern("Harmonic"), + VanityPattern { + prefix: "Harmonic".into(), + suffix: String::new(), + } + ); + } + + #[test] + fn prune_drops_dominated_patterns() { + let kept = prune_dominated(&[ + VanityPattern { + prefix: "Harmonic".into(), + suffix: "xyz".into(), + }, + VanityPattern { + prefix: "Harmonic".into(), + suffix: String::new(), + }, + VanityPattern { + prefix: "Cavey".into(), + suffix: "CooL".into(), + }, + ]); + assert_eq!( + kept, + vec![ + VanityPattern { + prefix: "Harmonic".into(), + suffix: String::new(), + }, + VanityPattern { + prefix: "Cavey".into(), + suffix: "CooL".into(), + }, + ] + ); + } } From 687667191f540277293838414d0ba39e7a19c215 Mon Sep 17 00:00:00 2001 From: cavey Date: Fri, 11 Sep 2026 18:12:04 +0000 Subject: [PATCH 2/2] add --count/--max-count and drop filled kinds from the matchers --- README.md | 20 +- kernels/base58.cu | 1 + kernels/base58.h | 3 + kernels/opencl/base58.cl | 9 + kernels/opencl/host.c | 14 + kernels/pattern_table.h | 6 +- kernels/vanity.cu | 7 + kernels/vanity.h | 1 + kernels/vanity_keypair.cu | 7 + kernels/vanity_keypair.h | 1 + src/fast/check_match.rs | 109 ++++--- src/fast/mod.rs | 85 ++--- src/main.rs | 635 ++++++++++++++++++++++++++++---------- 13 files changed, 643 insertions(+), 255 deletions(-) diff --git a/README.md b/README.md index d04321d..59c86da 100644 --- a/README.md +++ b/README.md @@ -40,7 +40,7 @@ If you don't have a GPU, consider using [vast.ai](https://cloud.vast.ai/?ref_id= ## Usage -vanity grinds three kinds of vanity output. All three share the same mining flags (`--num-cpus`, `--num-gpus`, `--count`) described in [Common options](#common-options). +vanity grinds three kinds of vanity output. All three share the same mining flags (`--num-cpus`, `--num-gpus`, `--count`, and for pattern grinds `--max-count`) described in [Common options](#common-options). | Command | Produces | How the address is derived | Speed | | --- | --- | --- | --- | @@ -60,8 +60,10 @@ controlled by `base` (the signer), not by a standalone keypair. You supply the `--base` pubkey (the signer for the create instruction), the program `--owner`, and one or more `--pattern`s to match (base58). A pubkey -hits if it matches **any** pattern. `...` splits prefix and suffix (`.` is -not in base58): +hits if it matches **any** pattern; the grind stops after `--count` hits of +**each** pattern. While a harder pattern is still short of `--count`, easier +kinds keep collecting up to `--max-count` (default: `--count`). `...` splits +prefix and suffix (`.` is not in base58): ```bash vanity grind \ @@ -70,7 +72,7 @@ vanity grind \ --pattern abc... \ --num-gpus 1 -# several targets in one run (OR): +# several targets in one run (continues until each has --count hits): vanity grind \ --base \ --owner BPFLoaderUpgradeab1e11111111111111111111111 \ @@ -114,7 +116,9 @@ pub fn create_account_with_seed( ### Grind a vanity keypair `grind-keypair` searches genuine ed25519 keypairs whose **public key** matches -any `--pattern` (same `Prefix...Suffix` syntax as `grind`). This is the classic +any `--pattern` (same `Prefix...Suffix` syntax as `grind`), and keeps going +until `--count` hits of each pattern (extras of an easy kind are kept up +to `--max-count`). This is the classic "vanity wallet/mint" case: the result is a standalone keypair you can sign with. It's slower than `grind` because each attempt computes an ed25519 scalar multiplication. The CPU path uses a batched custom ed25519 implementation @@ -186,9 +190,11 @@ doppler: 1/4 sign-extendable segment(s) | --- | --- | --- | --- | | `--num-cpus ` | all | `0` (= all logical cores) | CPU mining threads | | `--num-gpus ` | all (GPU builds only) | `1` | GPUs to mine on | -| `--count ` | all | `1` | stop after finding N matches | +| `--count ` | `grind`, `grind-keypair` | `1` | stop once **every** `--pattern` has N matches | +| `--max-count ` | `grind`, `grind-keypair` | `--count` | keep saving each `--pattern` up to N while waiting; then drop that kind from the search | +| `--count ` | `grind-doppler` | `1` | stop after N matching keypairs | | `--case-insensitive` | `grind`, `grind-keypair` | off | match patterns ignoring case | -| `--pattern` | `grind`, `grind-keypair` | — | repeatable `Prefix...Suffix` / `Prefix...` / `...Suffix`; OR'd | +| `--pattern` | `grind`, `grind-keypair` | — | repeatable `Prefix...Suffix` / `Prefix...` / `...Suffix`; a hit matches any unfilled kind | | `--segments <1-4>` | `grind-doppler` | `1` | sign-extendable segments required | To run purely on CPU (no GPU), build without a GPU feature, or pass `--num-gpus 0`. diff --git a/kernels/base58.cu b/kernels/base58.cu index becf975..508e73c 100644 --- a/kernels/base58.cu +++ b/kernels/base58.cu @@ -16,6 +16,7 @@ __constant__ uint8_t d_match_lut[58]; #include "pattern_table.h" __constant__ uint32_t d_n_patterns; +__constant__ unsigned long long d_active_mask = ~0ULL; __constant__ uint8_t d_prefix_len[VANITY_MAX_PATTERNS]; __constant__ uint8_t d_suffix_len[VANITY_MAX_PATTERNS]; __constant__ uint8_t d_prefix[VANITY_MAX_PATTERNS][VANITY_MAX_PATTERN_LEN]; diff --git a/kernels/base58.h b/kernels/base58.h index 4ea6bc0..edb28a5 100644 --- a/kernels/base58.h +++ b/kernels/base58.h @@ -16,6 +16,7 @@ extern __constant__ uint8_t d_match_lut[58]; #endif extern __constant__ uint32_t d_n_patterns; +extern __constant__ unsigned long long d_active_mask; extern __constant__ uint8_t d_prefix_len[VANITY_MAX_PATTERNS]; extern __constant__ uint8_t d_suffix_len[VANITY_MAX_PATTERNS]; extern __constant__ uint8_t d_prefix[VANITY_MAX_PATTERNS][VANITY_MAX_PATTERN_LEN]; @@ -189,6 +190,8 @@ static __device__ __forceinline__ bool fd_base58_check_match_any_32_words(const for (uint32_t pi = 0; pi < n_pat; pi++) { + if ((d_active_mask & (1ULL << pi)) == 0) + continue; uint8_t plen = d_prefix_len[pi]; uint8_t slen = d_suffix_len[pi]; diff --git a/kernels/opencl/base58.cl b/kernels/opencl/base58.cl index af7241e..c4cc787 100644 --- a/kernels/opencl/base58.cl +++ b/kernels/opencl/base58.cl @@ -14,6 +14,9 @@ #define VANITY_PT_SLEN (4 + VANITY_MAX_PATTERNS) #define VANITY_PT_PREF (VANITY_PT_SLEN + VANITY_MAX_PATTERNS) #define VANITY_PT_SUF (VANITY_PT_PREF + VANITY_MAX_PATTERNS * VANITY_MAX_PATTERN_LEN) +#define VANITY_PT_SUF_END (VANITY_PT_SUF + VANITY_MAX_PATTERNS * VANITY_MAX_PATTERN_LEN) +#define VANITY_PT_ACTIVE ((VANITY_PT_SUF_END + 7) & ~7) +#define VANITY_PT_SIZE (VANITY_PT_ACTIVE + 8) #endif __constant uchar base58_chars[] = "123456789ABCDEFGHJKLMNPQRSTUVWXYZabcdefghijkmnopqrstuvwxyz"; @@ -218,7 +221,13 @@ static bool fd_base58_check_match_any_32_words(const uint state[8], if (n_pat == 0) return true; + ulong active = 0UL; + for (uint b = 0; b < 8; b++) + active |= ((ulong)patterns[VANITY_PT_ACTIVE + b]) << (8 * b); + for (uint pi = 0; pi < n_pat; pi++) { + if ((active & (1UL << pi)) == 0) + continue; uchar plen = patterns[VANITY_PT_PLEN + pi]; uchar slen = patterns[VANITY_PT_SLEN + pi]; ulong pref = (ulong)VANITY_PT_PREF + (ulong)pi * (ulong)VANITY_MAX_PATTERN_LEN; diff --git a/kernels/opencl/host.c b/kernels/opencl/host.c index b2b9aca..acef645 100644 --- a/kernels/opencl/host.c +++ b/kernels/opencl/host.c @@ -334,6 +334,13 @@ void gpu_grind_read(void *opaque, uint8_t *out) { c->max_iters = adapt_iters(c->max_iters, elapsed, GRIND_ITERS_MIN, GRIND_ITERS_MAX); } +void gpu_grind_set_active_mask(void *opaque, unsigned long long mask) { + GrindCtx *c = (GrindCtx *)opaque; + CK(clEnqueueWriteBuffer(c->queue, c->patterns, CL_TRUE, + VANITY_PT_ACTIVE, sizeof(mask), &mask, + 0, NULL, NULL), "active mask"); +} + void gpu_grind_destroy(void *opaque) { GrindCtx *c = (GrindCtx *)opaque; clFinish(c->queue); @@ -483,6 +490,13 @@ void gpu_keypair_read(void *opaque, uint8_t *out) { c->max_iters = adapt_iters(c->max_iters, elapsed, KP_ITERS_MIN, KP_ITERS_MAX); } +void gpu_keypair_set_active_mask(void *opaque, unsigned long long mask) { + KeypairCtx *c = (KeypairCtx *)opaque; + CK(clEnqueueWriteBuffer(c->queue, c->patterns, CL_TRUE, + VANITY_PT_ACTIVE, sizeof(mask), &mask, + 0, NULL, NULL), "active mask"); +} + void gpu_keypair_destroy(void *opaque) { KeypairCtx *c = (KeypairCtx *)opaque; clFinish(c->queue); diff --git a/kernels/pattern_table.h b/kernels/pattern_table.h index 775f79a..92099f7 100644 --- a/kernels/pattern_table.h +++ b/kernels/pattern_table.h @@ -14,7 +14,9 @@ #define VANITY_PT_SLEN (4 + VANITY_MAX_PATTERNS) #define VANITY_PT_PREF (VANITY_PT_SLEN + VANITY_MAX_PATTERNS) #define VANITY_PT_SUF (VANITY_PT_PREF + VANITY_MAX_PATTERNS * VANITY_MAX_PATTERN_LEN) -#define VANITY_PT_SIZE (VANITY_PT_SUF + VANITY_MAX_PATTERNS * VANITY_MAX_PATTERN_LEN) +#define VANITY_PT_SUF_END (VANITY_PT_SUF + VANITY_MAX_PATTERNS * VANITY_MAX_PATTERN_LEN) +#define VANITY_PT_ACTIVE ((VANITY_PT_SUF_END + 7) & ~7) +#define VANITY_PT_SIZE (VANITY_PT_ACTIVE + 8) /* Packed blob from MatchTargets::gpu_blob → flat SoA. */ static void vanity_unpack_patterns(const uint8_t *blob, uint64_t blob_len, @@ -37,6 +39,8 @@ static void vanity_unpack_patterns(const uint8_t *blob, uint64_t blob_len, memcpy(out + VANITY_PT_SUF + i * VANITY_MAX_PATTERN_LEN, p, slen); p += slen; } + unsigned long long active = ~0ULL; + memcpy(out + VANITY_PT_ACTIVE, &active, 8); } #endif diff --git a/kernels/vanity.cu b/kernels/vanity.cu index 87d58e3..8075052 100644 --- a/kernels/vanity.cu +++ b/kernels/vanity.cu @@ -277,6 +277,13 @@ extern "C" void* gpu_grind_init( return (void *)ctx; } +extern "C" void gpu_grind_set_active_mask(void *opaque, unsigned long long mask) +{ + GpuGrindCtx *ctx = (GpuGrindCtx *)opaque; + cudaSetDevice(ctx->device_id); + cudaMemcpyToSymbol(d_active_mask, &mask, sizeof(mask)); +} + template __global__ void vanity_search(uint8_t *buffer, uint64_t stride, unsigned long long max_cycles); diff --git a/kernels/vanity.h b/kernels/vanity.h index 02a96d9..688fc12 100644 --- a/kernels/vanity.h +++ b/kernels/vanity.h @@ -9,5 +9,6 @@ extern "C" void gpu_grind_launch(void *ctx, uint8_t *seed); extern "C" int gpu_grind_query(void *ctx); extern "C" void gpu_grind_read(void *ctx, uint8_t *out); extern "C" void gpu_grind_destroy(void *ctx); +extern "C" void gpu_grind_set_active_mask(void *ctx, unsigned long long mask); #endif diff --git a/kernels/vanity_keypair.cu b/kernels/vanity_keypair.cu index 7553638..128c408 100644 --- a/kernels/vanity_keypair.cu +++ b/kernels/vanity_keypair.cu @@ -164,6 +164,13 @@ extern "C" void* gpu_keypair_init( return (void *)ctx; } +extern "C" void gpu_keypair_set_active_mask(void *opaque, unsigned long long mask) +{ + GpuKeypairCtx *ctx = (GpuKeypairCtx *)opaque; + cudaSetDevice(ctx->device_id); + cudaMemcpyToSymbol(d_active_mask, &mask, sizeof(mask)); +} + extern "C" void gpu_keypair_launch(void *opaque, uint8_t *seed) { GpuKeypairCtx *ctx = (GpuKeypairCtx *)opaque; diff --git a/kernels/vanity_keypair.h b/kernels/vanity_keypair.h index af0ead8..95700b1 100644 --- a/kernels/vanity_keypair.h +++ b/kernels/vanity_keypair.h @@ -8,6 +8,7 @@ extern "C" void gpu_keypair_launch(void *ctx, uint8_t *seed); extern "C" int gpu_keypair_query(void *ctx); extern "C" void gpu_keypair_read(void *ctx, uint8_t *out); extern "C" void gpu_keypair_destroy(void *ctx); +extern "C" void gpu_keypair_set_active_mask(void *ctx, unsigned long long mask); /* "doppler" keypair grind: matches ed25519 pubkeys with at least required_segments sign-extendable 32-bit segments (see kernels/opencl/ diff --git a/src/fast/check_match.rs b/src/fast/check_match.rs index faeb27c..22cc338 100644 --- a/src/fast/check_match.rs +++ b/src/fast/check_match.rs @@ -69,13 +69,14 @@ struct PatternIdx { /// Bytes are canonical raw_base58 indices, not ASCII. pub struct MatchTargets { patterns: Vec, - max_prefix_len: u8, - max_suffix_len: u8, match_lut: &'static [u8; 58], } impl MatchTargets { - pub fn new(patterns: &[(&str, &str)], case_insensitive: bool) -> Self { + pub fn new( + patterns: &[(&str, &str)], + case_insensitive: bool, + ) -> Self { assert!( patterns.len() <= MAX_PATTERNS, "at most {MAX_PATTERNS} patterns" @@ -88,8 +89,6 @@ impl MatchTargets { let lut = match_lut(case_insensitive); let mut out = Vec::with_capacity(patterns.len()); - let mut max_prefix_len = 0u8; - let mut max_suffix_len = 0u8; for &(prefix, suffix) in patterns { debug_assert!(prefix.len() <= MAX_PATTERN_LEN); debug_assert!(suffix.len() <= MAX_PATTERN_LEN); @@ -101,43 +100,65 @@ impl MatchTargets { for (i, &b) in suffix.as_bytes().iter().enumerate() { suffix_idx[i] = char_to_canonical(b, alphabet, lut); } - let plen = prefix.len() as u8; - let slen = suffix.len() as u8; - max_prefix_len = max_prefix_len.max(plen); - max_suffix_len = max_suffix_len.max(slen); out.push(PatternIdx { prefix_idx, - prefix_len: plen, + prefix_len: prefix.len() as u8, suffix_idx, - suffix_len: slen, + suffix_len: suffix.len() as u8, }); } Self { patterns: out, - max_prefix_len, - max_suffix_len, match_lut: lut, } } + /// Bit `i` is set when pattern `i` matches. Empty set → no match. + #[cfg(test)] + pub fn match_mask(&self, bytes: &[u8; 32]) -> u64 { + self.match_mask_active(bytes, u64::MAX) + } + + /// Bit `i` is set when pattern `i` matches and that kind is in `active`. #[inline] - pub fn matches(&self, bytes: &[u8; 32]) -> bool { - check_match_32(bytes, &self.patterns, self.match_lut) + pub fn match_mask_active( + &self, + bytes: &[u8; 32], + active: u64, + ) -> u64 { + check_match_32(bytes, &self.patterns, self.match_lut, active) } /// Canonical-index blob consumed by `gpu_grind_init` / `gpu_keypair_init`. + #[cfg(any(test, feature = "gpu"))] pub fn gpu_blob(&self) -> Vec { + let max_prefix_len = self + .patterns + .iter() + .map(|p| p.prefix_len) + .max() + .unwrap_or(0); + let max_suffix_len = self + .patterns + .iter() + .map(|p| p.suffix_len) + .max() + .unwrap_or(0); let mut v = Vec::with_capacity( 6 + self .patterns .iter() - .map(|p| 2 + p.prefix_len as usize + p.suffix_len as usize) + .map(|p| { + 2 + p.prefix_len as usize + p.suffix_len as usize + }) .sum::(), ); - v.extend_from_slice(&(self.patterns.len() as u32).to_le_bytes()); - v.push(self.max_prefix_len); - v.push(self.max_suffix_len); + v.extend_from_slice( + &(self.patterns.len() as u32).to_le_bytes(), + ); + v.push(max_prefix_len); + v.push(max_suffix_len); for p in &self.patterns { v.push(p.prefix_len); v.extend_from_slice(&p.prefix_idx[..p.prefix_len as usize]); @@ -194,7 +215,8 @@ fn check_match_32( bytes: &[u8; 32], patterns: &[PatternIdx], match_lut: &[u8; 58], -) -> bool { + active: u64, +) -> u64 { let mut in_leading_0s = 0usize; while in_leading_0s < 32 && bytes[in_leading_0s] == 0 { in_leading_0s += 1; @@ -247,12 +269,16 @@ fn check_match_32( let encoded_length = RAW58_SZ_32 - skip; if patterns.is_empty() { - return true; + return 1; } // Same per-char early-reject walk as the original single-pattern // matcher; extra patterns re-use already-emitted limbs. - for p in patterns { + let mut mask = 0u64; + for (pi, p) in patterns.iter().enumerate() { + if active & (1u64 << pi) == 0 { + continue; + } let mut ok = true; for i in 0..p.prefix_len as usize { let rb_idx = skip + i; @@ -288,10 +314,10 @@ fn check_match_32( } } if ok { - return true; + mask |= 1u64 << pi; } } - false + mask } #[cfg(test)] @@ -308,16 +334,16 @@ mod tests { ) -> bool { let s = fd_bs58::encode_32(*bytes); if ci { - let lc: String = s - .chars() - .map(|c| { - if c == 'L' { - c - } else { - c.to_ascii_lowercase() - } - }) - .collect(); + let lc: String = + s.chars() + .map(|c| { + if c == 'L' { + c + } else { + c.to_ascii_lowercase() + } + }) + .collect(); (prefix.is_empty() || lc.starts_with(prefix)) && (suffix.is_empty() || lc.ends_with(suffix)) } else { @@ -334,7 +360,7 @@ mod tests { let seed: [u8; 32] = h[..32].try_into().unwrap(); let bytes = pubkey_of_seed(&seed); assert_eq!( - target.matches(&bytes), + target.match_mask(&bytes) != 0, encode_matches(&bytes, "zz", "LM", false), "seed {i}" ); @@ -350,7 +376,7 @@ mod tests { let seed: [u8; 32] = h[..32].try_into().unwrap(); let bytes = pubkey_of_seed(&seed); assert_eq!( - target.matches(&bytes), + target.match_mask(&bytes) != 0, encode_matches(&bytes, "mithriL", "", true), "seed {i}" ); @@ -369,7 +395,7 @@ mod tests { let prefix = &key[..key.len().min(3)]; let suffix = &key[key.len().saturating_sub(2)..]; let target = MatchTargets::new(&[(prefix, suffix)], false); - assert!(target.matches(&bytes), "{key}"); + assert!(target.match_mask(&bytes) != 0, "{key}"); } } @@ -377,14 +403,15 @@ mod tests { fn or_of_two_patterns() { let key = "XkCriyrNwS3G4rzAXtG5B1nnvb5Ka1JtCku93VqeKAr"; let bytes = fd_bs58::decode_32(key).unwrap(); - let miss = MatchTargets::new(&[("zzz", ""), ("aaa", "qq")], false); - assert!(!miss.matches(&bytes)); + let miss = + MatchTargets::new(&[("zzz", ""), ("aaa", "qq")], false); + assert_eq!(miss.match_mask(&bytes), 0); let hit = MatchTargets::new(&[("zzz", ""), ("XkC", "KAr")], false); - assert!(hit.matches(&bytes)); + assert_eq!(hit.match_mask(&bytes), 0b10); let pref_only = MatchTargets::new(&[("nope", "nope"), ("XkC", "")], false); - assert!(pref_only.matches(&bytes)); + assert_eq!(pref_only.match_mask(&bytes), 0b10); } #[test] diff --git a/src/fast/mod.rs b/src/fast/mod.rs index aa72f02..0c79a97 100644 --- a/src/fast/mod.rs +++ b/src/fast/mod.rs @@ -9,7 +9,7 @@ use field::{batch_invert, Fe}; use group::{edwards_d2, Niels, Point}; use crate::{ - check_write_permissions, save_keypair, ABORTED, FOUND, + check_write_permissions, credit_kinds, save_keypair, ABORTED, TOTAL_ATTEMPTS, }; use sha2::{Digest, Sha512}; @@ -195,7 +195,11 @@ unsafe fn keygen_batch_simd( #[cfg(target_arch = "x86_64")] #[target_feature(enable = "avx512f,avx512ifma,avx512dq")] -unsafe fn grind_thread_simd(target: &MatchTargets, count: u32) { +unsafe fn grind_thread_simd( + target: &MatchTargets, + min_count: u32, + max_count: u32, +) { check_write_permissions(); let mut seeds: [[u8; 32]; BATCH] = [[0u8; 32]; BATCH]; for s in seeds.iter_mut() { @@ -206,7 +210,7 @@ unsafe fn grind_thread_simd(target: &MatchTargets, count: u32) { let mut local: u64 = 0; loop { - if grind_done(count) { + if grind_done(min_count) { break; } keygen_batch_simd(&mut seeds, &mut used, &mut pubkeys); @@ -216,21 +220,28 @@ unsafe fn grind_thread_simd(target: &MatchTargets, count: u32) { local = 0; } for j in 0..BATCH { - if target.matches(&pubkeys[j]) { - let prev = FOUND.fetch_add(1, Ordering::SeqCst); - if prev < count { - let s = fd_bs58::encode_32(pubkeys[j]); - eprintln!("\r\x1b[Kmatch: {s}"); + let mask = target.match_mask_active( + &pubkeys[j], + crate::unfilled_mask(max_count), + ); + if mask != 0 && credit_kinds(mask, max_count) { + let s = fd_bs58::encode_32(pubkeys[j]); + crate::ui_with_match(|| { + eprintln!("match: {s}"); eprintln!("pubkey: {s}"); save_keypair(&used[j], &pubkeys[j], &s); - } + }); } } } TOTAL_ATTEMPTS.fetch_add(local, Ordering::Relaxed); } -fn grind_thread_scalar(target: &MatchTargets, count: u32) { +fn grind_thread_scalar( + target: &MatchTargets, + min_count: u32, + max_count: u32, +) { check_write_permissions(); let mut seeds: [[u8; 32]; BATCH] = [[0u8; 32]; BATCH]; for s in seeds.iter_mut() { @@ -241,7 +252,7 @@ fn grind_thread_scalar(target: &MatchTargets, count: u32) { let mut local: u64 = 0; loop { - if grind_done(count) { + if grind_done(min_count) { break; } keygen_batch(&mut seeds, &mut used, &mut pubkeys); @@ -251,14 +262,17 @@ fn grind_thread_scalar(target: &MatchTargets, count: u32) { local = 0; } for j in 0..BATCH { - if target.matches(&pubkeys[j]) { - let prev = FOUND.fetch_add(1, Ordering::SeqCst); - if prev < count { - let s = fd_bs58::encode_32(pubkeys[j]); - eprintln!("\r\x1b[Kmatch: {s}"); + let mask = target.match_mask_active( + &pubkeys[j], + crate::unfilled_mask(max_count), + ); + if mask != 0 && credit_kinds(mask, max_count) { + let s = fd_bs58::encode_32(pubkeys[j]); + crate::ui_with_match(|| { + eprintln!("match: {s}"); eprintln!("pubkey: {s}"); save_keypair(&used[j], &pubkeys[j], &s); - } + }); } } } @@ -267,8 +281,7 @@ fn grind_thread_scalar(target: &MatchTargets, count: u32) { #[inline(always)] fn grind_done(count: u32) -> bool { - FOUND.load(Ordering::Relaxed) >= count - || ABORTED.load(Ordering::Relaxed) + crate::quotas_filled(count) || ABORTED.load(Ordering::Relaxed) } #[cfg(test)] @@ -285,22 +298,14 @@ pub fn pubkey_of_seed(seed: &[u8; 32]) -> [u8; 32] { use rayon::iter::{IntoParallelIterator, ParallelIterator}; -pub fn reset_grind() { - FOUND.store(0, Ordering::SeqCst); - TOTAL_ATTEMPTS.store(0, Ordering::SeqCst); - ABORTED.store(false, Ordering::SeqCst); +pub fn reset_grind(n_kinds: usize) { + crate::start_grind(n_kinds); } pub fn request_abort() { ABORTED.store(true, Ordering::SeqCst); } -/// Returns the previous found count (caller should print only if `prev < count`). -#[cfg(feature = "gpu")] -pub fn note_found() -> u32 { - FOUND.fetch_add(1, Ordering::SeqCst) -} - #[cfg(feature = "gpu")] pub fn is_done(count: u32) -> bool { grind_done(count) @@ -318,23 +323,31 @@ pub fn backend_name() -> &'static str { } } -/// Run batched CPU keypair workers until `count` matches or abort. -/// Call [`reset_grind`] first if coordinating with a GPU thread that shares -/// these counters via [`add_attempts`] / [`note_found`] / [`is_done`]. -pub fn run_cpu_workers(target: &MatchTargets, num_cpus: u32, count: u32) { +/// Run batched CPU keypair workers until `min_count` hits of each +/// pattern or abort. Extra hits are kept up to `max_count`. Call +/// [`reset_grind`] first if coordinating with a GPU thread that shares +/// these counters via [`is_done`]. +pub fn run_cpu_workers( + target: &MatchTargets, + num_cpus: u32, + min_count: u32, + max_count: u32, +) { (0..num_cpus) .into_par_iter() .for_each(|_| { #[cfg(target_arch = "x86_64")] { if simd::available() { - unsafe { grind_thread_simd(&target, count) }; + unsafe { + grind_thread_simd(target, min_count, max_count) + }; } else { - grind_thread_scalar(&target, count); + grind_thread_scalar(target, min_count, max_count); } } #[cfg(not(target_arch = "x86_64"))] - grind_thread_scalar(&target, count); + grind_thread_scalar(target, min_count, max_count); }); } diff --git a/src/main.rs b/src/main.rs index 5cca70a..1da45c5 100644 --- a/src/main.rs +++ b/src/main.rs @@ -30,8 +30,10 @@ use std::{ io::Write, str::FromStr, sync::{ - atomic::{AtomicBool, AtomicU32, AtomicU64, Ordering}, - Arc, + atomic::{ + AtomicBool, AtomicU32, AtomicU64, AtomicUsize, Ordering, + }, + Arc, Mutex, }, thread, time::{Duration, Instant}, @@ -48,6 +50,9 @@ pub enum Command { } /// Repeatable vanity target: a pubkey matches if it matches any `--pattern`. +/// Grind stops once every `--pattern` has `--count` hits. Extra hits of +/// an easy kind are kept up to `--max-count` (default: `--count`); that +/// kind is then dropped from the search. /// /// Syntax (`.` is not in base58, so `...` is an unambiguous split): /// `--pattern Cavey...CooL` prefix `Cavey` and suffix `CooL` @@ -56,7 +61,7 @@ pub enum Command { /// `--pattern Harmonic` prefix only (the `...` may be omitted) #[derive(Debug, Parser)] pub struct PatternArgs { - /// Vanity pattern. Repeatable; OR'd together. + /// Vanity pattern. Repeatable; a hit matches any unfilled kind. #[clap(long, action = clap::ArgAction::Append)] pub pattern: Vec, @@ -87,9 +92,14 @@ pub struct GrindArgs { #[clap(long, default_value_t = 0)] pub num_cpus: u32, - /// Number of matching addresses to find before stopping + /// Stop once every --pattern has this many matches #[clap(long, default_value_t = 1)] pub count: u32, + + /// Keep collecting each --pattern up to this many while waiting + /// for --count. Defaults to --count. + #[clap(long)] + pub max_count: Option, } #[derive(Debug, Parser)] @@ -106,9 +116,14 @@ pub struct GrindKeypairArgs { #[clap(long, default_value_t = 0)] pub num_cpus: u32, - /// Number of matching keypairs to find before stopping + /// Stop once every --pattern has this many matches #[clap(long, default_value_t = 1)] pub count: u32, + + /// Keep collecting each --pattern up to this many while waiting + /// for --count. Defaults to --count. + #[clap(long)] + pub max_count: Option, } #[derive(Debug, Parser)] @@ -188,7 +203,9 @@ pub struct DeployArgs { // ─── globals ──────────────────────────────────────────────────────────────── -pub(crate) static FOUND: AtomicU32 = AtomicU32::new(0); +pub(crate) static FOUND_PER: [AtomicU32; fast::MAX_PATTERNS] = + [const { AtomicU32::new(0) }; fast::MAX_PATTERNS]; +pub(crate) static N_KINDS: AtomicUsize = AtomicUsize::new(1); pub(crate) static TOTAL_ATTEMPTS: AtomicU64 = AtomicU64::new(0); /// Set by the Ctrl-C handler; makes every grind loop wind down. Needed /// because with no explicit handler, SIGINT relies on the kernel's default @@ -196,9 +213,93 @@ pub(crate) static TOTAL_ATTEMPTS: AtomicU64 = AtomicU64::new(0); /// container (common on GPU cloud hosts), so Ctrl-C would otherwise be ignored. pub(crate) static ABORTED: AtomicBool = AtomicBool::new(false); +pub(crate) fn reset_quotas(n_kinds: usize) { + assert!((1..=fast::MAX_PATTERNS).contains(&n_kinds), "kind count"); + N_KINDS.store(n_kinds, Ordering::SeqCst); + for slot in FOUND_PER.iter() { + slot.store(0, Ordering::SeqCst); + } +} + +/// Bits for kinds that still have room under `quota` (at most `quota` each). +pub(crate) fn unfilled_mask(quota: u32) -> u64 { + let n = N_KINDS.load(Ordering::Relaxed); + let mut mask = 0u64; + for i in 0..n { + if FOUND_PER[i].load(Ordering::Relaxed) < quota { + mask |= 1u64 << i; + } + } + mask +} + +/// Credit every *unfilled* kind in `mask`. Caps each kind at `quota`. +/// Returns true if any kind was still under quota (print / save this hit). +pub(crate) fn credit_kinds(mask: u64, quota: u32) -> bool { + let mask = mask & unfilled_mask(quota); + if mask == 0 { + return false; + } + let n = N_KINDS.load(Ordering::Relaxed); + let mut useful = false; + for i in 0..n { + if mask & (1u64 << i) == 0 { + continue; + } + loop { + let prev = FOUND_PER[i].load(Ordering::SeqCst); + if prev >= quota { + break; + } + if FOUND_PER[i] + .compare_exchange( + prev, + prev + 1, + Ordering::SeqCst, + Ordering::SeqCst, + ) + .is_ok() + { + useful = true; + break; + } + } + } + useful +} + +pub(crate) fn quotas_filled(quota: u32) -> bool { + let n = N_KINDS.load(Ordering::Relaxed); + (0..n).all(|i| FOUND_PER[i].load(Ordering::Relaxed) >= quota) +} + +pub(crate) fn start_grind(n_kinds: usize) { + TOTAL_ATTEMPTS.store(0, Ordering::SeqCst); + ABORTED.store(false, Ordering::SeqCst); + reset_quotas(n_kinds); +} + fn done(target: u32) -> bool { - FOUND.load(Ordering::SeqCst) >= target - || ABORTED.load(Ordering::SeqCst) + quotas_filled(target) || ABORTED.load(Ordering::SeqCst) +} + +/// `--count` is the stop threshold; `--max-count` is the per-kind save cap. +fn resolve_counts(count: u32, max_count: Option) -> (u32, u32) { + let count = count.max(1); + let max_count = max_count.unwrap_or(count); + assert!( + max_count >= count, + "--max-count ({max_count}) must be >= --count ({count})" + ); + (count, max_count) +} + +fn targets_header(count: u32, max_count: u32) -> String { + if max_count == count { + format!("targets ({count} of each):") + } else { + format!("targets ({count} of each, max {max_count}):") + } } // ─── bs58 probability (from cavemanloverboy/bs58p) ────────────────────────── @@ -360,45 +461,48 @@ impl VanitySpec { .collect() } - fn matches(&self, pubkey: &str) -> bool { - self.patterns.iter().any(|p| { - matches_target( + fn match_mask(&self, pubkey: &str) -> u64 { + let mut mask = 0u64; + for (i, p) in self.patterns.iter().enumerate() { + if matches_target( pubkey, &p.prefix, &p.suffix, self.case_insensitive, - ) - }) - } - - fn label(&self) -> String { - self.patterns - .iter() - .map(|p| format_target_label(&p.prefix, &p.suffix)) - .collect::>() - .join(" | ") - } - - fn probability(&self) -> f64 { - self.patterns - .iter() - .map(|p| { - bs58_probability( - &p.prefix, - &p.suffix, - self.case_insensitive, - ) - }) - .sum::() - .min(1.0) + ) { + mask |= 1u64 << i; + } + } + mask } - fn expected_attempts(&self) -> f64 { - let p = self.probability(); - if p <= 0.0 { - f64::INFINITY - } else { - 1.0 / p + fn live_targets(&self, count: u32, max_count: u32) -> LiveTargets { + LiveTargets { + count, + max_count, + rows: self + .patterns + .iter() + .map(|p| { + let probability = bs58_probability( + &p.prefix, + &p.suffix, + self.case_insensitive, + ); + let expected = if probability <= 0.0 { + f64::INFINITY + } else { + count.max(1) as f64 / probability + }; + TargetRow { + label: format_target_label( + &p.prefix, &p.suffix, + ), + probability, + expected, + } + }) + .collect(), } } } @@ -428,29 +532,9 @@ fn parse_pattern(raw: &str) -> VanityPattern { VanityPattern { prefix, suffix } } -/// `q` matches every string that `p` matches (and possibly more). -fn dominates(q: &VanityPattern, p: &VanityPattern) -> bool { - p.prefix.starts_with(&q.prefix) && p.suffix.ends_with(&q.suffix) -} - -fn prune_dominated(patterns: &[VanityPattern]) -> Vec { - let mut out = Vec::new(); - for p in patterns { - if out.iter().any(|q| dominates(q, p)) { - continue; - } - out.retain(|q| !dominates(p, q)); - out.push(p.clone()); - } - out -} - fn resolve_spec(args: &PatternArgs) -> VanitySpec { let raw = &args.pattern; - assert!( - !raw.is_empty(), - "supply at least one --pattern" - ); + assert!(!raw.is_empty(), "supply at least one --pattern"); assert!( raw.len() <= fast::MAX_PATTERNS, "at most {} --pattern flags", @@ -458,7 +542,7 @@ fn resolve_spec(args: &PatternArgs) -> VanitySpec { ); let mut patterns = Vec::with_capacity(raw.len()); - for r in &raw { + for r in raw { let pat = parse_pattern(r); validate_bs58("pattern prefix", &pat.prefix); validate_bs58("pattern suffix", &pat.suffix); @@ -474,13 +558,37 @@ fn resolve_spec(args: &PatternArgs) -> VanitySpec { }); } VanitySpec { - patterns: prune_dominated(&patterns), + patterns, case_insensitive: args.case_insensitive, } } // ─── formatting ───────────────────────────────────────────────────────────── +struct TargetRow { + label: String, + probability: f64, + expected: f64, +} + +struct LiveTargets { + count: u32, + max_count: u32, + rows: Vec, +} + +struct UiState { + live: Option, + start: Option, + painted: bool, +} + +static UI: Mutex = Mutex::new(UiState { + live: None, + start: None, + painted: false, +}); + fn format_duration(secs: f64) -> String { if secs < 0.0 { return "any moment".into(); @@ -498,6 +606,172 @@ fn format_duration(secs: f64) -> String { format!("{}d {}h", s / 86400, (s % 86400) / 3600) } +fn format_expected_attempts(expected: f64) -> String { + if expected.is_finite() { + (expected as u64).to_formatted_string(&Locale::en) + } else { + "∞".into() + } +} + +fn format_eta( + remaining_hits: u32, + probability: f64, + rate: f64, +) -> String { + if remaining_hits == 0 { + return "done".into(); + } + if rate <= 0.0 || probability <= 0.0 { + return "—".into(); + } + format_duration((remaining_hits as f64 / probability) / rate) +} + +fn format_kind_eta( + found: u32, + count: u32, + max_count: u32, + probability: f64, + rate: f64, +) -> String { + if found >= max_count { + return "capped".into(); + } + if found >= count { + return "collecting".into(); + } + format_eta(count.saturating_sub(found), probability, rate) +} + +fn format_target_live_line( + row: &TargetRow, + found: u32, + count: u32, + max_count: u32, + rate: f64, +) -> String { + format!( + " {} {}/{} | {:.6e} | expected {} | ETA {}", + row.label, + found, + max_count, + row.probability, + format_expected_attempts(row.expected), + format_kind_eta(found, count, max_count, row.probability, rate), + ) +} + +fn format_stats_line(total: u64, rate: f64, elapsed: f64) -> String { + format!( + "{} attempts | {} attempts/sec | elapsed: {}", + total.to_formatted_string(&Locale::en), + (rate as u64).to_formatted_string(&Locale::en), + format_duration(elapsed), + ) +} + +fn ui_rates(start: Instant) -> (u64, f64, f64) { + let total = TOTAL_ATTEMPTS.load(Ordering::Relaxed); + let elapsed = start.elapsed().as_secs_f64().max(1e-9); + (total, total as f64 / elapsed, elapsed) +} + +/// Cursor stays on the stats line (no trailing newline) after a paint. +fn paint_live(st: &mut UiState, first: bool) { + let Some(ref live) = st.live else { + return; + }; + let Some(start) = st.start else { + return; + }; + let (total, rate, elapsed) = ui_rates(start); + let t = live.rows.len(); + if !first && st.painted && t > 0 { + eprint!("\x1b[{t}A"); + } + for (i, row) in live.rows.iter().enumerate() { + let found = FOUND_PER[i].load(Ordering::Relaxed); + eprint!( + "\r\x1b[K{}\n", + format_target_live_line( + row, + found, + live.count, + live.max_count, + rate, + ) + ); + } + eprint!("\r\x1b[K{}", format_stats_line(total, rate, elapsed)); + let _ = std::io::stderr().flush(); + st.painted = true; +} + +fn clear_live(st: &mut UiState) { + if !st.painted { + return; + } + let t = st + .live + .as_ref() + .map(|l| l.rows.len()) + .unwrap_or(0); + if t > 0 { + eprint!("\x1b[{t}A"); + } + eprint!("\r\x1b[J"); + let _ = std::io::stderr().flush(); + st.painted = false; +} + +fn ui_start(live: LiveTargets, start: Instant) { + let mut st = UI.lock().unwrap(); + st.live = Some(live); + st.start = Some(start); + st.painted = false; + paint_live(&mut st, true); +} + +fn ui_refresh() { + let mut st = UI.lock().unwrap(); + if st.live.is_none() { + return; + } + paint_live(&mut st, false); +} + +pub(crate) fn ui_with_match(f: impl FnOnce()) { + let mut st = UI.lock().unwrap(); + clear_live(&mut st); + f(); + paint_live(&mut st, true); +} + +fn ui_stop() { + let mut st = UI.lock().unwrap(); + if st.painted { + eprintln!(); + st.painted = false; + } + st.live = None; +} + +fn spawn_ui_reporter( + shutdown: Arc, +) -> thread::JoinHandle<()> { + thread::spawn(move || loop { + if shutdown.load(Ordering::SeqCst) { + break; + } + thread::sleep(Duration::from_secs(1)); + if shutdown.load(Ordering::SeqCst) { + break; + } + ui_refresh(); + }) +} + fn print_status(total: u64, rate: f64, elapsed: f64, expected: f64) { let e_time = if rate > 0.0 && expected.is_finite() { format!( @@ -548,6 +822,8 @@ struct GpuOps { query: unsafe extern "C" fn(*mut c_void) -> i32, read: unsafe extern "C" fn(*mut c_void, *mut u8), destroy: unsafe extern "C" fn(*mut c_void), + /// Drop filled kinds from the device matcher. None for doppler. + set_active_mask: Option, } /// Drive `num_gpus` async kernel launches until the grind is done. @@ -562,10 +838,20 @@ fn run_gpu_workers( init: impl Fn(u32) -> *mut c_void, done: impl Fn() -> bool, mut on_result: impl FnMut(usize, &[u8], f64), + quota: u32, ) { let contexts: Vec<*mut c_void> = (0..num_gpus).map(&init).collect(); let mut launch_times = vec![Instant::now(); num_gpus as usize]; let mut in_flight = vec![false; num_gpus as usize]; + let push_mask = |ctx: *mut c_void| { + let Some(set) = ops.set_active_mask else { + return; + }; + unsafe { set(ctx, unfilled_mask(quota)) }; + }; + for &ctx in &contexts { + push_mask(ctx); + } let read = |ctx: *mut c_void| -> ([u8; OUT], u64) { let mut out = [0u8; OUT]; @@ -601,6 +887,7 @@ fn run_gpu_workers( in_flight[i] = false; if !done() { + push_mask(ctx); let seed = new_gpu_seed(i as u32); launch_times[i] = Instant::now(); unsafe { (ops.launch)(ctx, seed.as_ptr()) }; @@ -779,21 +1066,19 @@ pub fn deploy_with_max_program_len_with_seed( fn grind(mut args: GrindArgs) { maybe_update_num_cpus(&mut args.num_cpus); let spec = resolve_spec(&args.spec); - let expected = spec.expected_attempts(); - let prob = spec.probability(); + let (target_count, max_count) = + resolve_counts(args.count, args.max_count); + start_grind(spec.patterns.len()); + let live = spec.live_targets(target_count, max_count); #[cfg(feature = "gpu")] eprintln!("using {} cpus, {} gpus", args.num_cpus, args.num_gpus); #[cfg(not(feature = "gpu"))] eprintln!("using {} cpus", args.num_cpus); - eprintln!( - "target: {} | probability: {:.6e} | expected: {} attempts", - spec.label(), - prob, - (expected as u64).to_formatted_string(&Locale::en) - ); + eprintln!("{}", targets_header(target_count, max_count)); - let target_count = args.count; let shutdown = Arc::new(AtomicBool::new(false)); + let grind_start = Instant::now(); + ui_start(live, grind_start); let spec = Arc::new(spec); #[cfg(feature = "gpu")] @@ -818,6 +1103,9 @@ fn grind(mut args: GrindArgs) { query: gpu_grind_query, read: gpu_grind_read, destroy: gpu_grind_destroy, + set_active_mask: Some( + gpu_grind_set_active_mask, + ), }, |id| unsafe { gpu_grind_init( @@ -837,20 +1125,29 @@ fn grind(mut args: GrindArgs) { .chain_update(owner) .finalize() .into(); - let out_str = fd_bs58::encode_32(reconstructed); + let out_str = + fd_bs58::encode_32(reconstructed); - if spec.matches(&out_str) { - eprintln!( - "\r\x1b[Kgpu {} match: {} in {:.3}s", - i, out_str, time_sec - ); - eprintln!( - "out seed = {out:?} -> {}", - core::str::from_utf8(&out[..16]).unwrap() - ); - FOUND.fetch_add(1, Ordering::SeqCst); + let mask = spec.match_mask(&out_str); + if mask != 0 + && credit_kinds(mask, max_count) + { + ui_with_match(|| { + eprintln!( + "gpu {} match: {} in {:.3}s", + i, out_str, time_sec + ); + eprintln!( + "out seed = {out:?} -> {}", + core::str::from_utf8( + &out[..16] + ) + .unwrap() + ); + }); } }, + max_count, ) }) .unwrap(), @@ -859,12 +1156,7 @@ fn grind(mut args: GrindArgs) { None }; - let grind_start = Instant::now(); - let reporter = spawn_hashrate_reporter( - Arc::clone(&shutdown), - expected, - grind_start, - ); + let reporter = spawn_ui_reporter(Arc::clone(&shutdown)); (0..args.num_cpus).into_par_iter().for_each(|i| { let timer = Instant::now(); @@ -900,22 +1192,20 @@ fn grind(mut args: GrindArgs) { local_batch -= 4096; } - if spec.matches(&pubkey) { + let mask = spec.match_mask(&pubkey); + if mask != 0 && credit_kinds(mask, max_count) { if local_batch > 0 { TOTAL_ATTEMPTS.fetch_add(local_batch, Ordering::Relaxed); local_batch = 0; } let time_secs = timer.elapsed().as_secs_f64(); - let elapsed_global = grind_start.elapsed().as_secs_f64().max(1e-9); - let total = TOTAL_ATTEMPTS.load(Ordering::Relaxed); - let global_rate = total as f64 / elapsed_global; - eprintln!( - "\r\x1b[Kcpu {i} match: {pubkey}; {seed:?} -> {} in {:.3}s; {} attempts/sec", - core::str::from_utf8(&seed).unwrap(), - time_secs, - (global_rate as u64).to_formatted_string(&Locale::en) - ); - FOUND.fetch_add(1, Ordering::SeqCst); + ui_with_match(|| { + eprintln!( + "cpu {i} match: {pubkey}; {seed:?} -> {} in {:.3}s", + core::str::from_utf8(&seed).unwrap(), + time_secs, + ); + }); if done(target_count) { break; } @@ -930,6 +1220,7 @@ fn grind(mut args: GrindArgs) { shutdown.store(true, Ordering::SeqCst); reporter.join().unwrap(); + ui_stop(); let total = TOTAL_ATTEMPTS.load(Ordering::Relaxed); let elapsed = grind_start @@ -938,7 +1229,7 @@ fn grind(mut args: GrindArgs) { .max(1e-9); let rate = total as f64 / elapsed; eprintln!( - "\r\x1b[Kdone: {} attempts in {} at {} attempts/sec", + "done: {} attempts in {} at {} attempts/sec", total.to_formatted_string(&Locale::en), format_duration(elapsed), (rate as u64).to_formatted_string(&Locale::en) @@ -951,8 +1242,10 @@ fn grind_keypair(mut args: GrindKeypairArgs) { check_write_permissions(); maybe_update_num_cpus(&mut args.num_cpus); let spec = resolve_spec(&args.spec); - let expected = spec.expected_attempts(); - let prob = spec.probability(); + let (target_count, max_count) = + resolve_counts(args.count, args.max_count); + fast::reset_grind(spec.patterns.len()); + let live = spec.live_targets(target_count, max_count); #[cfg(feature = "gpu")] eprintln!( "using {} cpus, {} gpus (cpu backend: {})", @@ -966,41 +1259,20 @@ fn grind_keypair(mut args: GrindKeypairArgs) { args.num_cpus, fast::backend_name() ); - eprintln!( - "target: {} | probability: {:.6e} | expected: {} attempts", - spec.label(), - prob, - (expected as u64).to_formatted_string(&Locale::en) - ); - - let target_count = args.count; - fast::reset_grind(); + eprintln!("{}", targets_header(target_count, max_count)); let shutdown = Arc::new(AtomicBool::new(false)); let grind_start = Instant::now(); + ui_start(live, grind_start); let match_targets = { let pairs = spec.pairs(); fast::MatchTargets::new(&pairs, spec.case_insensitive) }; #[cfg(feature = "gpu")] let spec = Arc::new(spec); + #[cfg(not(feature = "gpu"))] + let _ = spec; - // Reporter reads the shared fast-path counters (CPU + GPU both update them). - let reporter = { - let shutdown = Arc::clone(&shutdown); - thread::spawn(move || loop { - if shutdown.load(Ordering::SeqCst) { - break; - } - thread::sleep(Duration::from_secs(1)); - if shutdown.load(Ordering::SeqCst) { - break; - } - let elapsed = grind_start.elapsed().as_secs_f64(); - let total = fast::total_attempts(); - let rate = total as f64 / elapsed.max(1e-9); - print_status(total, rate, elapsed, expected); - }) - }; + let reporter = spawn_ui_reporter(Arc::clone(&shutdown)); #[cfg(feature = "gpu")] let gpu_thread = if args.num_gpus > 0 { @@ -1018,6 +1290,9 @@ fn grind_keypair(mut args: GrindKeypairArgs) { query: gpu_keypair_query, read: gpu_keypair_read, destroy: gpu_keypair_destroy, + set_active_mask: Some( + gpu_keypair_set_active_mask, + ), }, |id| unsafe { gpu_keypair_init( @@ -1038,13 +1313,16 @@ fn grind_keypair(mut args: GrindKeypairArgs) { SigningKey::from_bytes(&found_seed); let pubkey_bytes = signing_key.verifying_key().to_bytes(); - let pubkey_str = fd_bs58::encode_32(pubkey_bytes); + let pubkey_str = + fd_bs58::encode_32(pubkey_bytes); - if spec.matches(&pubkey_str) { - let prev = fast::note_found(); - if prev < target_count { + let mask = spec.match_mask(&pubkey_str); + if mask != 0 + && credit_kinds(mask, max_count) + { + ui_with_match(|| { eprintln!( - "\r\x1b[Kgpu {} match: {} in {:.3}s", + "gpu {} match: {} in {:.3}s", i, pubkey_str, time_sec ); eprintln!("pubkey: {pubkey_str}"); @@ -1053,9 +1331,10 @@ fn grind_keypair(mut args: GrindKeypairArgs) { &pubkey_bytes, &pubkey_str, ); - } + }); } }, + max_count, ) }) .unwrap(), @@ -1064,7 +1343,12 @@ fn grind_keypair(mut args: GrindKeypairArgs) { None }; - fast::run_cpu_workers(&match_targets, args.num_cpus, target_count); + fast::run_cpu_workers( + &match_targets, + args.num_cpus, + target_count, + max_count, + ); // CPU workers finished (found enough or aborted); stop GPU too. fast::request_abort(); @@ -1076,6 +1360,7 @@ fn grind_keypair(mut args: GrindKeypairArgs) { shutdown.store(true, Ordering::SeqCst); reporter.join().unwrap(); + ui_stop(); let total = fast::total_attempts(); let elapsed = grind_start @@ -1084,7 +1369,7 @@ fn grind_keypair(mut args: GrindKeypairArgs) { .max(1e-9); let rate = total as f64 / elapsed; eprintln!( - "\r\x1b[Kdone: {} attempts in {} at {} attempts/sec", + "done: {} attempts in {} at {} attempts/sec", total.to_formatted_string(&Locale::en), format_duration(elapsed), (rate as u64).to_formatted_string(&Locale::en) @@ -1101,9 +1386,11 @@ fn grind_doppler(mut args: DopplerArgs) { "--segments must be between 1 and 4" ); + let target_count = args.count; + start_grind(1); let prob = doppler_probability(args.segments); let expected = if prob > 0.0 { - 1.0 / prob + target_count as f64 / prob } else { f64::INFINITY }; @@ -1119,7 +1406,6 @@ fn grind_doppler(mut args: DopplerArgs) { (expected as u64).to_formatted_string(&Locale::en) ); - let target_count = args.count; let shutdown = Arc::new(AtomicBool::new(false)); #[cfg(feature = "gpu")] @@ -1137,6 +1423,7 @@ fn grind_doppler(mut args: DopplerArgs) { query: gpu_doppler_query, read: gpu_doppler_read, destroy: gpu_doppler_destroy, + set_active_mask: None, }, |id| unsafe { gpu_doppler_init(id as i32, segments) }, || done(target_count), @@ -1148,7 +1435,9 @@ fn grind_doppler(mut args: DopplerArgs) { let pubkey_bytes = signing_key.verifying_key().to_bytes(); - if doppler_count_segments(&pubkey_bytes) >= segments + if doppler_count_segments(&pubkey_bytes) + >= segments + && credit_kinds(1, target_count) { let pubkey_str = fd_bs58::encode_32(pubkey_bytes); @@ -1156,11 +1445,18 @@ fn grind_doppler(mut args: DopplerArgs) { "\r\x1b[Kgpu {} match: {} in {:.3}s", i, pubkey_str, time_sec ); - print_doppler_result(&found_seed, &pubkey_str); - save_keypair(&found_seed, &pubkey_bytes, &pubkey_str); - FOUND.fetch_add(1, Ordering::SeqCst); + print_doppler_result( + &found_seed, + &pubkey_str, + ); + save_keypair( + &found_seed, + &pubkey_bytes, + &pubkey_str, + ); } }, + target_count, ) }) .unwrap(), @@ -1199,7 +1495,9 @@ fn grind_doppler(mut args: DopplerArgs) { local_batch -= 4096; } - if doppler_count_segments(&pubkey_bytes) >= segments { + if doppler_count_segments(&pubkey_bytes) >= segments + && credit_kinds(1, target_count) + { if local_batch > 0 { TOTAL_ATTEMPTS.fetch_add(local_batch, Ordering::Relaxed); local_batch = 0; @@ -1216,7 +1514,6 @@ fn grind_doppler(mut args: DopplerArgs) { ); print_doppler_result(&seed, &pubkey_str); save_keypair(&seed, &pubkey_bytes, &pubkey_str); - FOUND.fetch_add(1, Ordering::SeqCst); if done(target_count) { break; } @@ -1373,7 +1670,7 @@ pub(crate) fn save_keypair( .open(&path) .and_then(|mut f| f.write_all(json.as_bytes())) { - Ok(_) => println!("keypair generated at: ./{}", path), + Ok(_) => eprintln!("keypair generated at: ./{}", path), Err(err) => { eprintln!("failed to write keypair to {path}: {err}") } @@ -1475,6 +1772,10 @@ extern "C" { pub fn gpu_grind_query(ctx: *mut std::ffi::c_void) -> i32; pub fn gpu_grind_read(ctx: *mut std::ffi::c_void, out: *mut u8); pub fn gpu_grind_destroy(ctx: *mut std::ffi::c_void); + pub fn gpu_grind_set_active_mask( + ctx: *mut std::ffi::c_void, + mask: u64, + ); pub fn gpu_keypair_init( id: i32, @@ -1489,6 +1790,10 @@ extern "C" { pub fn gpu_keypair_query(ctx: *mut std::ffi::c_void) -> i32; pub fn gpu_keypair_read(ctx: *mut std::ffi::c_void, out: *mut u8); pub fn gpu_keypair_destroy(ctx: *mut std::ffi::c_void); + pub fn gpu_keypair_set_active_mask( + ctx: *mut std::ffi::c_void, + mask: u64, + ); pub fn gpu_doppler_init( id: i32, @@ -1577,33 +1882,23 @@ mod tests { } #[test] - fn prune_drops_dominated_patterns() { - let kept = prune_dominated(&[ - VanityPattern { - prefix: "Harmonic".into(), - suffix: "xyz".into(), - }, - VanityPattern { - prefix: "Harmonic".into(), - suffix: String::new(), - }, - VanityPattern { - prefix: "Cavey".into(), - suffix: "CooL".into(), - }, - ]); - assert_eq!( - kept, - vec![ + fn match_mask_sets_bit_per_pattern() { + let spec = VanitySpec { + patterns: vec![ VanityPattern { - prefix: "Harmonic".into(), + prefix: "XkC".into(), suffix: String::new(), }, VanityPattern { - prefix: "Cavey".into(), - suffix: "CooL".into(), + prefix: "zzz".into(), + suffix: String::new(), }, - ] - ); + ], + case_insensitive: false, + }; + let key = "XkCriyrNwS3G4rzAXtG5B1nnvb5Ka1JtCku93VqeKAr"; + assert_eq!(spec.match_mask(key), 0b01); + assert_eq!(spec.match_mask("zzzABC"), 0b10); + assert_eq!(spec.match_mask("nope"), 0); } }