NVIDIA GPU acceleration - #121
Open
levinster82 wants to merge 26 commits into
Open
Conversation
- main.c: fix borrow detection in savecheckpoint() and carry detection
in checkpoint load; both had wrong multi-byte arithmetic that could
silently corrupt the deterministic seed offset on resume, causing
passphrase-mode keys to be re-generated from already-searched space
- main.c: reject negative -n arguments instead of wrapping to SIZE_MAX
via size_t cast, which caused the tool to run indefinitely
- worker.c: fix shiftpk() out-of-bounds read of src[PUBLIC_LEN] on the
last loop iteration when src points to a 32-byte pk_batch entry; also
fix shift-by-8 UB when sbits==0 after the modulo
- worker.c: check writetofile() return values in onionready() and error
out rather than silently producing incomplete key directories
- worker.c/worker.h: change endwork from volatile int to
volatile sig_atomic_t as required by C11 for signal-handler access
- ioutil.c: fix five do{}while(0) EINTR retry loops where continue
exited the loop instead of retrying; most critically closefile() was
returning 0 (success) on EINTR
- yaml.c: check writetofile() return values in yamlin_parseandcreate()
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Implements optional NVIDIA GPU support (--enable-cuda) that replaces the CPU thread pool for single-word filter searches. Key design points: - Each GPU thread runs an independent copy of the worker_batch loop with Montgomery batch inversion (BATCHNUM_GPU points per batch) to amortize the dominant fe_invert cost (~231 fe_mul) down to ~4-5 fe_mul/point - Buffer layout uses [step][thread_id] ordering for coalesced warp access - ge_eightpoint and filter table in __constant__ memory (broadcast cache) - Persistent kernel (no relaunch overhead); stops via __device__ cuda_endwork - BATCHNUM selected at runtime from VRAM, dispatched via template<int B> with 5 explicit instantiations (64/128/256/512/1024) - Result ring uses mapped pinned memory + __threadfence_system() + per-slot done flags to eliminate cudaMemcpy from the hot drain path - GPU parameters (blocks, threads/block, batchnum) auto-tuned at runtime via cudaGetDeviceProperties; build-time -gencode flags from configure probe - Falls back to CPU for PASSPHRASE mode and numwords > 1 New files: ed25519/cuda/fe_cuda.cuh, ed25519/cuda/ge_cuda.cuh, keccak_cuda.cuh, worker_cuda.h, gpu_autoconf.cu, worker_cuda.cu, GPU_BUILD_PLAN.txt, GPU_TODO.txt Modified: configure.ac, GNUmakefile.in, main.c, worker.c, worker.h Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
- Add shiftpk_cuda device function (direct port of worker.c shiftpk): shifts a 32-byte public key left by N bits, safe for in-place use - Add multi-word check loop in kernel: after initial filter match, shift pk by filter_bits and check next word, matching CPU numwords behavior - Fix upload_filters() to populate filter_bits for each gpu_filter_entry: INTFILTER uses popcount(mask); BINFILTER uses flen*8 + leading mask bits - Fix OMITMASK INTFILTER case: was reading .m from a struct that lacks it; now correctly reads from the global ifiltermask - Remove numwords > 1 CPU fallback from main.c (no longer needed) Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
- Kernel spins on done[slot]==0 before writing to a reserved slot, ensuring the drain thread has consumed the previous result there - Spin also exits on cuda_endwork to prevent deadlock if a stop signal arrives while the ring is full - Ring size increased from 1024 to 4096 to reduce spin frequency under high match rates (e.g. short 1-3 char filters) Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Six fixes needed to make the CUDA build link and run on hardware: - gpu_autoconf.cu, worker_cuda.cu: wrap project C headers in extern "C" to prevent C++ name mangling of symbols compiled as C objects; add missing vec.h include before worker.h (VEC_STRUCT dependency) - gpu_autoconf.cu: replace hardcoded ref10 symbol names and phantom ge_eightpoint with CRYPTO_NAMESPACE() calls and a new ge_get_base_precomp() accessor; fixes link failure when building with donna or other impls - ed25519/ref10/ge.h, ge_scalarmult_base.c: add ge_get_base_precomp() to expose base[0][0] (generator B in Duif form) for GPU constant init - GNUmakefile.in: use nvcc as linker when ENABLE_CUDA=1 so libcudart is resolved without requiring it in the system library path - types.h: add #include <stdint.h> for uint8_t etc. in C++ compilation - worker_cuda.cu: remove #define _POSIX_C_SOURCE (conflicts with CUDA headers which define a newer version) Tested on RTX 3070 (sm_86): GPU detected, addresses generated correctly. Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Adds calc/sec and succ/sec reporting for GPU mode, matching the output format of the CPU -s flag and honouring the same reportdelay/realtimestats settings. Implementation: - gpu_state: h_numcalc/d_numcalc mapped pinned counter; CPU reads directly, GPU writes via atomicAdd (one per block per outer loop, not per thread) - kernel_args: endwork_flag and numcalc passed as device pointers instead of __device__ globals, eliminating cudaMemcpyToSymbol stream serialization - drain_thread: periodic stats reporting; local_success tracks matched keys independently of keysgenerated (which is only incremented under -n) - gpu_worker_launch: accepts reportdelay and realtimestats from main() Also fixes clean exit after -n N: cudaMemcpyToSymbol(cuda_endwork) was serialized behind the persistent kernel causing a deadlock; now the drain thread writes *h_endwork = 1 directly to mapped pinned memory. Tested on RTX 3070: ~850M calc/sec, succ/sec correct for 6-char prefix. Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Always initialize istarttime at drain thread startup (not gated on reportdelay) so elapsed time is tracked unconditionally. After the drain loop exits, emit one cumulative stats line to stderr using total h_numcalc and local_success over the full run duration. This fires on both clean exit (-n N) and Ctrl+C, giving throughput and elapsed time for every run regardless of whether -s was passed. Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
- Link libcudart statically (-cudart static) so binary runs on any machine without requiring CUDA runtime installed - configure auto-detects nvcc in PATH and /usr/local/cuda/bin; --enable-cuda now defaults to auto instead of no - Binary falls back to CPU threads silently when no GPU is found; prints "no GPU found, using N CPU threads" as a single combined line - Add -C flag to force CPU-only mode in CUDA builds - Print final stats line on CPU exit (matching GPU behavior) - Fix sk scalar offset: step is 1 not 8 (cuda_ge_eightpoint = 1*B) - Silence scary CUDA errors on no-GPU machines via cudaGetDeviceCount Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
fe_invert_cuda had two bugs in its addition chain: 1. Missing one squaring in the early steps (z^4 instead of z^8) 2. t0=z^11 was overwritten in place, so the final multiply used z^31 instead of z^11, producing z^(2^255-1) not z^(2^255-21) This caused every GPU-generated public key to be cryptographically wrong while the stored scalar was correct, so sk->pk always failed verification. pk->hostname passed because the hostname was computed from the same (wrong) pk bytes inside the kernel. Fix: rewrite fe_invert_cuda as a direct port of ref10/pow225521.h, preserving t0=z^11 untouched through the full chain exactly as the reference implementation does. Also commit previously-staged security fixes to gpu_autoconf.cu: sodium_memzero for h_sk after upload and h_results before free. Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Verifies a mkp224o-generated key directory without running Tor:
1. sk->pk: derives public key from scalar using pure Python ed25519
(no dependencies); cross-checks with PyNaCl if available
2. pk->hostname: rederives .onion address from public key per v3 spec
Used to confirm the fe_invert_cuda fix: GPU keys now pass both checks.
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Replace raw >calc/sec:..., succ/sec:..., rest/sec:... with a line that shows elapsed, speed, match ratio, ETA 50%/90%, and found count: > elapsed: 5s | speed: 855.1M/s | 1:34.4B | ETA 50%: 27s 90%: 1m32s | found: 0 Filter difficulty uses exact probability sum (1/Σ(1/d_i)) for BINFILTER, popcount(ifiltermask)/n for INTFILTER. No -lm dependency. GPU and CPU paths share print_stats_line() via statline.h. Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
- Add NVIDIA CUDA Toolkit as optional build dependency with install instructions for Debian/Ubuntu and Fedora/RHEL - Add GPU Acceleration section documenting auto-detection, -C flag, and RTX 3070 benchmark (850 M keys/s, ~34x over 16-core CPU) - Update -s flag description to show new human-readable stats format with ETA 50%/90% example output - Add prefix length difficulty/ETA table anchored to RTX 3070 numbers - Add [CUDA] reference link Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
The GPU batch kernel was stepping by 1*G and adding 1 per step to the scalar. Because valid ed25519 scalars require sk[0] & 7 == 0 (the cofactor-8 clamp), only every 8th filter match passed the sanity check, silently discarding 7/8 of hits. The ETA was consequently 8× too optimistic, and a 9-char search took ~3.87 days to 50% probability instead of the displayed ~8 hours. Fix: add ge_get_eightpoint_precomp() to ref10 (computes 8*G in Z=1 precomp form by calling ge_scalarmult_base(8), normalising with fe_invert, and extracting xy2d from ge_p3_to_cached), use it in gpu_init instead of base[0][0] (= 1*G), and multiply the scalar offset by 8 so the key pair remains consistent. With 8*G steps every produced scalar is a multiple of 8, so the sanity check passes unconditionally and the displayed ETA is now accurate. Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Commit 1df3d2e added ge_get_eightpoint_precomp() to ref10 with an extern declaration in ge.h but omitted the CRYPTO_NAMESPACE #define that every ref10 symbol needs. The definition therefore compiled as the bare name while gpu_autoconf.cu referenced the namespaced crypto_sign_ed25519_donna_ge_get_eightpoint_precomp, producing an undefined-reference link error in CUDA builds. Add the missing #define so the definition picks up the active namespace. Verified: compiling ge_scalarmult_base.c with -DED25519_donna now emits crypto_sign_ed25519_donna_ge_get_eightpoint_precomp. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
The "found:" stat showed keysgenerated, which onionready() only incremented inside the numneedgenerate (-n) block. Without -n it stayed 0 for the whole run even as keys were found and written to disk, which is confusing. Lift the increment out so it always counts; keep the early-out and endwork stop gated on -n so quota behavior is unchanged. Covers both CPU workers and the GPU drain thread, which share onionready(). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Reflects 0c2370e — keysgenerated now counts every found key regardless of the -n quota. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
A single process uses one CUDA device; explain running one process per card to use them all, including mismatched models (each process auto-configures to its own card and runs independently). Throughput scales linearly; effective rate is the sum of per-process speeds. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
The CUDA worker is a persistent kernel (one launch looping until the CPU sets endwork), which hangs Nsight Compute since it replays and awaits the kernel. Add an opt-in MKP_PROFILE_ITERS env knob: when set, the kernel exits after N candidates/thread and the launch releases the drain thread so the app exits cleanly. Default (unset) is unchanged — runs forever. PROFILING.md documents the two-phase plan: nsys for a quick bound/occupancy read on the normal binary, ncu (application replay, under the cap) for the detailed section analysis, plus a decision tree mapping the numbers to which optimization to pursue (memory layout vs int64-free field core vs occupancy). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
--gpu-metrics-device was renamed to --gpu-metrics-devices; document the ERR_NVGPUCTRPERM fix (sudo, or NVreg_RestrictProfilingToAdminUsers=0). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
The batch buffer stored X,Y,Z (30 int32/candidate) and read them back after batch inversion; X was used only for the public key's x-sign bit, which lives in byte 31 and never lands in a vanity prefix. Store only Y,Z (20 int32) and filter on the sign-less key. On a match the GPU emits just the secret scalar; the CPU drain thread recomputes the exact key (correct sign), SHA3 checksum and onion address — cheap since matches are rare. This cuts ~1/3 of the dominant DRAM traffic and pulls SHA3/formatting out of the hot kernel. Measured on an RTX 3070: ~843 -> ~995 M/s (+18%), DRAM throughput 82.5% -> 72%. Keys verified end-to-end (sk->pk, pk->hostname). Smaller slots also let the auto-tuner fit batchnum=1024. Also adds an MKP_BATCHNUM override (with a VRAM-budget guard) for tuning experiments. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
The batch inversion wrote every inverted Z back to global memory and the second pass read it back to convert/filter. Split the inversion into a forward-only helper (prefix products + acc = (prod Z)^-1) and run the backward pass inline in the kernel: each Z^-1 is derived and consumed (tobytes + filter) in registers, so the inverted Z never round-trips through global memory. acc is advanced before any early-out so candidates that fail the filter still keep the chain correct. Removes ~1/3 of the remaining batch-buffer traffic. On an RTX 3070: ~995 M/s -> ~1.04 G/s, DRAM 72% -> 61% (kernel is now compute/memory balanced rather than memory-bound). No register spills. Keys verified end-to-end (sk->pk, pk->hostname). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Document the ~1.23x result (843 M/s -> ~1.04 G/s on a 3070), the two structural wins, and the measured dead ends — notably why the int64->PTX field core was declined (kernel is latency-bound at 16.7% occupancy, not throughput-bound) and why occupancy is register-capped. Saves re-walking the path; flags that conclusions should be re-profiled on Ada/Blackwell. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
The drop-X and fused-inversion changes raised the RTX 3070 from ~850 M/s to ~1.04 G/s (~40x vs a 16-core CPU), measured and verified. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Two profile-guided structural wins, both correctness-verified: - drop X from the batch buffer; CPU recomputes the exact key on a hit - fuse the batch-inversion backward pass with the filter check 843 M/s -> ~1.04 G/s; DRAM 82.5% -> 61% (kernel now compute/memory balanced). Includes the profiling kit/hook and findings in PROFILING.md. No CPU-path code changed. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
This branch adds optional GPU support (NVIDIA/CUDA) to speed up the vanity address search, hitting around 850 million keys/sec on an RTX 3070 — far faster than CPU. It turns on automatically if you have a GPU and the CUDA tools installed, needs no special flags, and falls back to CPU on its own when there's no GPU (or use -C to force CPU).