lumbda/examples/cuda-fanout/plans/form-A-secp256k1-batched-mul.md

Form A — cuda-secp256k1-batched-mul implementation plan

Status: design only. No binary built. No remote GPU touched.

Reference catalog entry: ../CATALOG.md § A. Target: a secp256k1-batch-mul daemon binary that speaks the BSHK binary wire protocol the worker already drives for shake256-fanout, sitting next to it in examples/cuda-fanout/.

1. Upstream selection — pick FixedPaul/VanitySearch-Bitcrack

You would fork FixedPaul/VanitySearch-Bitcrack, not JeanLucPons. Reasons:

• License. FixedPaul ships AGPL-3.0 (https://github.com/FixedPaul/VanitySearch-Bitcrack), matching shake256-fanout.cu's AGPLv3 header. JeanLucPons is GPL-3.0 (https://github.com/JeanLucPons/VanitySearch). AGPL ↔ AGPL stays a clean derivative.
• Activity. FixedPaul last touched master 2025-05-16; JeanLucPons last commit 2022-04-04 (commits pages on each repo). Three years stale vs actively maintained.
• Speed. FixedPaul publishes 6.9 Gkeys/s on 4090, 8.8 Gkeys/s on 5090 (repo README). Their "optimized CUDA modular math" patch over upstream is the win we want.

No 3090 number published; you would measure on 3090-ai.foxhop.net once a binary exists. Expect ~3 Gkeys/s extrapolating from 3080's 2.65 Gkeys/s in our catalog.

2. Kernel extraction — files you need

Upstream layout (https://github.com/JeanLucPons/VanitySearch/tree/master/GPU; fork preserves shape):

• GPU/GPUMath.h (~900 lines) — __device__ modular arithmetic with inline PTX: _ModMult, _ModSqr, ModSub256, _ModInv, and crucially _ModInvGrouped (Montgomery batched inversion: one inverse per GRP_SIZE/2+1 points). Keep verbatim.
• GPU/GPUGroup.h (~1024 lines) — precomputed Gx[][], Gy[][] for G, 2G, … 512G plus _2Gnx/_2Gny. Windowed-mul ladder. Keep verbatim.
• GPU/GPUCompute.h (~450 lines) — ComputeKeys() Shamir-trick loop (P+iG and P-iG share one delta-x). Keep field arithmetic; strip SHA256 + RIPEMD160 + base58 + address-prefix branches.
• GPU/GPUEngine.cu (~920 lines, 5 __global__ kernels: comp_keys, comp_keys_p2sh, comp_keys_comp, comp_keys_pattern, comp_keys_p2sh_pattern) — drop entirely. Write one new secp_mul_batch_kernel(scalars[N], base, points[N]) calling _ModInvGrouped.

Drop: main.cpp, Vanity.cpp, Base58.cpp, Bech32.cpp, Wildcard.cpp, hash/* (SSE host hashes), GPUBase58.h, GPUHash.h, GPUWildcard.h, GPUGenerate.cpp — 19 .cpp files per upstream Makefile, zero kept.

License: copy each file's header verbatim, add SPDX line, drop a LICENSE-vanitysearch-bitcrack beside our existing license note.

3. Wire protocol — BSHK request, BSHR response

Mirror the existing shake protocol (shake256-fanout.cu lines 286-292):

Request, little-endian throughout, no alignment padding:

magic       u8[4]   = "BSHK"
op_id       u32     = 0x01 (secp_mul_batch)
n           u32     = number of scalars
base_x      u8[32]  base point X (big-endian field element, secp256k1 native)
base_y      u8[32]  base point Y
scalars     u8[n*32]  big-endian 256-bit ints

Response:

magic       u8[4]   = "BSHR" on success, "BERR" on failure
status      u32     = 0 ok, nonzero error code (1 bad-magic, 2 bad-n,
                     3 cuda-launch-fail, 4 truncated-input)
n           u32     = echoed scalar count
points      u8[n*64]  N points, each (X32 || Y32) big-endian; affine,
                     infinity encoded as all-zero

Endianness: LE u32 headers (matches shake binary at shake256-fanout.cu:286), BE field elements (matches secp256k1 native; ships straight to openssl/python-ecdsa for validation). Worker's wire-send-raw treats payload as opaque (gpu-worker.lsp:264), no change there.

4. gpu-worker.lsp handler sketch

You add one handler analogous to handle-binary-shake at gpu-worker.lsp:253. The daemon stays warm for the whole worker lifetime — same pattern, same daemon-process helper:

(define *binary-secp-mul* "./secp256k1-batch-mul")

;; in main: register at boot, paying CUDA init once.
(register-daemon! 'cuda-secp-mul-batch *binary-secp-mul*)

;; new binary dispatcher: payload[0..4] = "BSPK" magic to disambiguate
;; from "BSHK" shake requests.
(define (handle-binary-secp-mul client payload)
  (let* ((daemon  (cdr (assoc 'cuda-secp-mul-batch *daemons*)))
         (in-path  (gensym-path "/tmp/bend-secp-in"  ".bin"))
         (out-path (gensym-path "/tmp/bend-secp-out" ".bin"))
         (blob (substring payload 4 (string-length payload))))
    (write-binary-file in-path blob)
    (let ((status (daemon-process daemon in-path out-path #t)))
      (cond
        ((eq? status 'ok)
         (let ((result (read-binary-file out-path)))
           (delete-file in-path) (delete-file out-path)
           (wire-send-raw client (string-append "BSPR" result))))
        (else
         (delete-file in-path) (delete-file out-path)
         (wire-send-raw client (string-append "BERR" (cdr status))))))))

Extend handle-one (gpu-worker.lsp:271) to branch on first 4 bytes: BSHK → shake, BSPK → secp-mul. S-expression path stays for ping and any text-mode ops. CUDA context init runs twice at startup, never per request.

5. Daemon protocol — match shake256-fanout exactly

secp256k1-batch-mul --daemon reads stdin one cmd per line, emits one response per line, identical to shake256-fanout.cu:226-269:

ready                      ; once at startup after cudaFree(0)
process-bin <in> <out>     ; binary request
done <out>                 ; success
error <reason>             ; failure
quit / bye                 ; shutdown

daemon-process in gpu-worker.lsp:71 already writes process-bin <in> <out>\n & expects done/error — zero handler-side change. start-daemon (line 62) waits for ready before serving, so CUDA init blocks worker readiness deterministically.

CLI fallback: secp256k1-batch-mul --binary <in> <out> for one-shot test (matches shake256-fanout.cu:272).

6. Build recipe

Two options:

• A — submodule. git submodule add https://github.com/FixedPaul/VanitySearch-Bitcrack third_party/vanitysearch-bitcrack. Pro: trivial upstream-bump. Con: ~12 MB of unused .cpp / .vcxproj noise.
• B — vendor three headers. Copy GPUMath.h, GPUGroup.h, GPUCompute.h into examples/cuda-fanout/vendor/secp256k1-vanitysearch/ with AGPL-3.0 + SPDX preserved, add a PROVENANCE.md carrying source commit SHA. Pro: lean, audit-friendly. Con: manual sync.

Pick B — ~2400 lines, mirrors how shake256-fanout.cu carries its tiny-sha3 distillation inline.

Makefile additions:

NVCCFLAGS_SECP ?= -O3 -arch=sm_86 -Xcompiler="-O3 -Wall" \
                  -maxrregcount=0 --ptxas-options=-v

all: shake256-fanout secp256k1-batch-mul

secp256k1-batch-mul: secp256k1-batch-mul.cu \
                     vendor/secp256k1-vanitysearch/GPUMath.h \
                     vendor/secp256k1-vanitysearch/GPUGroup.h \
                     vendor/secp256k1-vanitysearch/GPUCompute.h
	$(NVCC) $(NVCCFLAGS_SECP) -Ivendor/secp256k1-vanitysearch -o $@ $<

test-secp: secp256k1-batch-mul test_secp_roundtrip.py
	./secp256k1-batch-mul --binary /tmp/secp-in.bin /tmp/secp-out.bin
	python3 test_secp_roundtrip.py /tmp/secp-in.bin /tmp/secp-out.bin

bench-secp: secp256k1-batch-mul bench_secp.py
	python3 bench_secp.py ./secp256k1-batch-mul

sm_86 covers 3090; sm_89 covers 4090. Build a fat binary with -gencode arch=compute_86,code=sm_86 -gencode arch=compute_89,code=sm_89 for both fleet hosts.

Worker auto-registration: one line, (register-daemon! 'cuda-secp-mul-batch *binary-secp-mul*) next to the shake registration at gpu-worker.lsp:299.

7. Validation plan

Correctness. test_secp_roundtrip.py runs N=100 known scalars (1, 2, 3, …, plus edge cases: curve order n → infinity, n-1, scalars near 2^128, 2^256-1). Computes k*G on CPU via coincurve (or ecdsa.SECP256k1.generator * k), byte-compares GPU output. Halt on first mismatch. Cross-check at N=10 against openssl EC_POINT_mul to catch a coincurve defect masking a kernel defect.

Performance. bench_secp.py sweeps N ∈ {1k, 10k, 100k, 1M}, records wire_ms, kernel_ms (extend response to carry u32 kernel_us after n), gpu/cpu ratio. CPU baseline: in-process coincurve loop. Median over 20 runs. Format mirrors bench_tiers.py:174 so the summary table renders unchanged.

Expected: 3090 ~3 Gkeys/s, 4090 ~6.9 Gkeys/s. CPU baseline ~50 kkeys/s single-thread coincurve → ~60,000× per core, ~5,000× vs 12-thread. Below 1,000× signals a defect — most likely PCIe transfer saturating (1M × 64 B = 64 MB out, ~2 ms at PCIe 4.0).

8. Risks & open questions

• License audit. FixedPaul declares AGPL-3.0 at repo level but GPU/*.h headers may still carry GPLv3 inherited from JeanLucPons. Audit each vendored file before commit. If a file is GPLv3-only, our binary becomes GPLv3 (still compatible with our AGPLv3 worker via §7 of GPLv3).
• Endomorphism in fork. FixedPaul claims "optimized CUDA modular math" but lists no specifics. If the speedup is a private GPUMath.h patch, vendoring upstream headers alone loses it. Diff FixedPaul/GPU/GPUMath.h vs JeanLucPons/GPU/GPUMath.h first; take FixedPaul's.
• Host RAM at N=1M. 1M × 64 B = 64 MB out, 1M × 32 B = 32 MB in. Worker holds ~3 copies (in-portal + daemon + out-portal): peak ~300 MB transient. Bounded on 24 GB / 80 GB hosts.
• GRP_SIZE preservation. Upstream's GRP_SIZE=1024 means each thread does 1024 keys sequentially sharing one _ModInvGrouped call. A naive one-thread-per-scalar kernel loses the amortisation — ~50× worse. New kernel launches ceil(N/1024) threads, each handling a 1024-scalar group; N rounds up to multiple of 1024 with sentinel scalar = 1, pad outputs discarded.
• Device memory. 1M-point request allocates ~96 MB device-side (in/out + scratch). cudaFree between requests to avoid fragmentation over 10k+ requests. Add explicit free in process_one_bin.

9. Effort estimate

Experienced CUDA dev, no fleet-side work counted:

step	days
vendor headers, audit licenses, write PROVENANCE.md	0.5
write secp_mul_batch_kernel reusing _ModInvGrouped group loop	1.5
daemon shell (stdin loop, --binary, --daemon modes) — copy shake	0.5
binary I/O, BSHK/BSHR parse + emit, host-side scalar layout	0.5
Makefile, smoke test on local nvcc (host has no GPU — compile only)	0.5
test_secp_roundtrip.py + coincurve oracle + openssl cross-check	1.0
bench_secp.py + sweeps + first 3090/4090 measurement	1.0
gpu-worker.lsp handler + dispatch branch + integration test	0.5
catalog promotion (live row, measured speedup)	0.0 (fox's call)
total	6 days

Concentration: ~40% kernel-extraction (steps 1+2), ~25% validation, ~15% bench harness, ~10% daemon/wire, ~10% handler. Kernel-extraction risk dominates; everything else mirrors shake256-fanout.

URLs cited

• Catalog § A: ../CATALOG.md
• Existing shake daemon: ../shake256-fanout.cu
• Worker handler pattern: ../gpu-worker.lsp lines 62, 71, 253, 264, 299
• Wire bench reference: ../bench_tiers.py lines 150-158, 174
• VanitySearch original: https://github.com/JeanLucPons/VanitySearch
• VanitySearch GPU folder: https://github.com/JeanLucPons/VanitySearch/tree/master/GPU
• VanitySearch last commit: https://github.com/JeanLucPons/VanitySearch/commits/master
• Bitcrack fork: https://github.com/FixedPaul/VanitySearch-Bitcrack
• Bitcrack last commit: https://github.com/FixedPaul/VanitySearch-Bitcrack/commits/master
• gECC paper (speedup reference): https://arxiv.org/pdf/2501.03245