This is the last post in a three-part build log about mining Neptune Cash. The first built a TIP5 accelerator on eight AMD Virtex UltraScale+ FPGAs on AWS EC2 F2. The second brought it onto an on-desk Osprey E300 — three VU35P UltraScale+ FPGAs with a Zynq 7010 as controller-only — running the whole miner in programmable logic with no soft processor. Both were good UltraScale+ hardware. Neither is what I mine with today.
By October I had spent the better part of six months in the FPGA development loop. Then I rewrote the whole thing in CUDA and, on a single NVIDIA L40S, hit the throughput target I had been chasing — in about one month. This post is about why that happened, and why it wasn’t really a surprise. Code in ai-hpc/tip5-neptune under gpu_pow/.
Table of contents
Open Table of contents
The six-month wall
Nothing was wrong with the FPGA designs. The S-box parallelization hit a real 98.4× latency win; the PL-only Zynq build taught me more about hardware than anything I’d done. The problem was the loop.
Every idea cost a full synthesis + place-and-route + timing-closure cycle to test. Minutes on the small Zynq, most of an hour on a VU35P, longer for an F2 AFI build. A wrong guess about pipeline depth or DSP packing wasn’t a thirty-second recompile — it was a coffee break, and sometimes a DRC error at the end that sent you back to the start. Six months in, the bottleneck was not the hardware’s hash rate. It was my iteration rate. And the thing about mining a live network is that the target keeps moving while you iterate.
So I ran an experiment I expected to lose: give the whole problem to a single GPU and see how far I get in a month.
Why TIP5 is secretly a GPU workload
The thing that made the GPU experiment win wasn’t raw FLOPS. It was the field.
TIP5 is defined over the Goldilocks field, p = 2^64 − 2^32 + 1. That prime is not arbitrary — it is chosen so that modular reduction is almost free. Because 2^64 ≡ 2^32 − 1 (mod p), a 128-bit product folds back into 64 bits with a couple of 32-bit shifts, an add, and a conditional subtraction — no 128-bit division, no Barrett or Montgomery table lookup strictly required. A GPU has a native 64-bit integer multiply and runs tens of thousands of threads; the Goldilocks reduction is a handful of ALU ops that every one of those threads can do independently.
That is the whole insight. On an FPGA you spend weeks building DSP multiplier trees and fighting timing to get a parallel field multiplier. On a GPU, the field multiply is already the hardware’s native operation, and the parallelism you fought for in fabric is just the thread grid. The mining loop becomes embarrassingly simple:
// one thread = one nonce; launch millions
__global__ void tip5_guess(const uint64_t* job_digest,
uint64_t nonce_base,
uint64_t threshold,
uint64_t* winner) {
uint64_t nonce = nonce_base + blockIdx.x * blockDim.x + threadIdx.x;
uint64_t state[16];
load_state(state, job_digest, nonce);
#pragma unroll
for (int r = 0; r < NUM_ROUNDS; ++r) { // 5 rounds
sbox_layer(state); // 4× split-and-lookup + 12× x^7
mds_layer(state); // 16×16 MDS, Goldilocks mul
add_round_constants(state, r);
}
if (digest_below(state, threshold))
atomicMin(winner, nonce);
}No PCIe round-trip per hash, no DMA batching gymnastics, no AXI master to hand-build — the nonce space is the launch grid. The gpu_pow/ tree is essentially a march toward making this kernel faster: tip5_hash_cuda_optimized.cu, then tip5_hash_cuda_building_blocks_optimized.cu, then tip5_hash_cuda_ultra_optimized.cu, with a tip5_hash_target_60.cu that names the goal — 60 M hashes/second — in the filename. A single L40S reached that target. Eight cloud FPGAs were built to chase it.
The NVIDIA L40S is, on paper, a graphics-and-inference card — Ada Lovelace, 48 GB GDDR6, 864 GB/s, 18,176 CUDA cores — not a mining ASIC. It didn’t need to be. Once the field arithmetic is native and the parallelism is free, “fast enough” arrives quickly, and the engineering moves from making one hash fast to keeping the SMs saturated: occupancy, register pressure per thread (16 field elements of state is a lot of registers), unrolling the 5 rounds, and laying out the MDS multiply so it stays in registers instead of spilling.
Mining is the guesser — proving is a different project
It’s worth being precise about what “mining” is here, because it’s easy to overclaim. A Neptune miner runs the proof-of-work guesser: the tip5_guess loop above, hunting for a nonce whose TIP5 digest clears the threshold. That is what this repo and this post are about, and it is genuinely well-suited to a GPU — millions of independent attempts, no shared state, no data dependency between nonces.
Generating the zk-STARK proofs that validate Neptune’s transactions is a completely separate — and much larger — piece of work. TIP5 is the hash inside those proofs, but proof generation is a different beast entirely: number-theoretic transforms over the Goldilocks field, low-degree extension, Merkle commitments, FRI folding, constraint evaluation. That’s a dedicated zk-STARK proof generator (the Triton VM prover), not a miner, and it dwarfs the guesser in scope. I treat it as its own project, and I’m not going to pretend the miner grew into it. The only thing the two share is the field: once you’ve written fast Goldilocks arithmetic in CUDA for one, the other becomes approachable — which is exactly the bridge I describe at the end.
The honest comparison
I want to be fair to the FPGAs, because the comparison people expect — raw hashes-per-watt — is not the one that decided it.
| UltraScale+ FPGA | Single L40S GPU | |
|---|---|---|
| Edit→test loop | synth + P&R + timing: minutes to an hour+ | nvcc recompile: seconds |
| Time to hit the target | ~6 months, still tuning | ~1 month |
| Unit to reach the goal | up to 8× VU47P (F2 instance) | one card |
| Field arithmetic | hand-built DSP multiplier trees, timing-closed | native 64-bit multiply, free across threads |
| Scaling out | replicate the AFI across 8 FPGAs | widen the launch grid, add cards |
| Where it genuinely wins | perf/watt at fixed function, huge scale | iteration speed, time-to-correct, flexibility |
At massive scale and fixed function, FPGAs (and ASICs beyond them) still win on energy per hash — that’s not in dispute. But I am one engineer iterating against a live network, and for that situation the GPU’s seconds-long compile-run loop and the Goldilocks field’s native fit on CUDA were decisive. The FPGA path optimized the hash; the GPU path let me optimize the whole miner fast enough to keep up with a moving target.
Where this led
Today the miner runs on the L40S, and the FPGA designs are a well-documented detour. But the deeper thing the GPU opened up wasn’t the miner itself — it was everything adjacent to it. Writing fast Goldilocks-field arithmetic in CUDA, profiling it under Nsight, and learning where a 64-bit-modular kernel actually spends its time is the skill that carried into the project I’m proudest of: a dedicated zk-STARK proof generator — a separate and far larger effort than this miner — that I built around register blocking and PTX-level field math, eventually 245 hand-written CUDA kernels on a Blackwell RTX PRO 6000. That’s the Blackwell post: a different project, not a sequel to this one. The shared DNA is only the field arithmetic.
The arc across these three posts is, in retrospect, the arc of how I actually think now. Start where the hardware instinct points (FPGA, fixed function, parallel fabric). Hit the wall that the instinct hides (iteration speed, time-to-result). Find the place where the math and the machine already agree (Goldilocks on CUDA), and pour everything into making that fast. That last step — squeezing a GPU harder than anyone else on a problem that genuinely fits it — is the work I want to be doing for the rest of my career.