cavemanloverboy · Sep 11, 2026 · Sep 10, 2026 · Sep 11, 2026
diff --git 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 |
 | --- | --- | --- | --- |
@@ -59,13 +59,25 @@ 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; 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 \
   --base <YOUR_WALLET_PUBKEY> \
   --owner BPFLoaderUpgradeab1e11111111111111111111111 \
-  --prefix abc \
+  --pattern abc... \
+  --num-gpus 1
+
+# several targets in one run (continues until each has --count hits):
+vanity grind \
+  --base <YOUR_WALLET_PUBKEY> \
+  --owner BPFLoaderUpgradeab1e11111111111111111111111 \
+  --pattern Cavey...CooL \
+  --pattern Harmonic... \
   --num-gpus 1
 ```
 
@@ -103,15 +115,18 @@ 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`), 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
+(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
@@ -175,9 +190,11 @@ doppler: 1/4 sign-extendable segment(s)
 | --- | --- | --- | --- |
 | `--num-cpus <N>` | all | `0` (= all logical cores) | CPU mining threads |
 | `--num-gpus <N>` | all (GPU builds only) | `1` | GPUs to mine on |
-| `--count <N>` | 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 |
+| `--count <N>` | `grind`, `grind-keypair` | `1` | stop once **every** `--pattern` has N matches |
+| `--max-count <N>` | `grind`, `grind-keypair` | `--count` | keep saving each `--pattern` up to N while waiting; then drop that kind from the search |
+| `--count <N>` | `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`; 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/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
@@ -1,3 +1,5 @@
+#include <string.h>
+#include <cuda_runtime.h>
 #include "base58.h"
 
 __device__ uint8_t const base58_chars[] = "123456789ABCDEFGHJKLMNPQRSTUVWXYZabcdefghijkmnopqrstuvwxyz";
@@ -11,6 +13,30 @@ __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__ 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];
+__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
@@ -10,6 +10,23 @@ 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__ 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];
+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 +129,105 @@ 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++)
+    {
+        if ((d_active_mask & (1ULL << pi)) == 0)
+            continue;
+        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
@@ -6,6 +6,19 @@
        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)
+#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";
 /* ci = case insensitive: map all uppercase to lowercase except for L. */
 __constant uchar base58_chars_ci[] = "123456789abcdefghjkLmnpqrstuvwxyzabcdefghijkmnopqrstuvwxyz";
@@ -79,8 +92,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 +164,101 @@ 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;
+
+    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;
+        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
+}