From 3a8c66a1480aaa022d20f2f9193464d4fb728cc7 Mon Sep 17 00:00:00 2001 From: "russell@unturf.com" Date: Fri, 5 Jun 2026 19:07:00 -0400 Subject: [PATCH] examples/cuda-fanout: Day-4 v4 stack landed (byte-clean, -12% vs v3); daemon flips to v3 --window-w 4 MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Day-4 task per fox: stack v3 windowed-G ladder (Day-3) with v2 Montgomery batch inversion (Day-2). Idea: Day-3 cut scalar_mul, residual ModInv now matters — which is what Day-2 needed to win. Result: - v4 (--window-w 4 --batch-inv) lands byte-identical vs coincurve at n in {32, 1000, 10000, 100000, 1000000}. - 3090-ai best-of-3 @ n=1M: v1 7.86 / v3 13.83 / v4 12.16 Mkeys/s. - v4 regresses -12% vs v3 because v2's Phase B/D walks run one thread per block (3906 active threads at n=1M; 3090 has ~125k concurrent thread budget). v3's per-thread ModInv saturates 1M parallel threads on 82 SMs — the threading model beats the smaller field-mult count. - HARD-RULE triggered: v4 < 1.10x v3 → no v4 promotion. Daemon flip per task brief option A: - start-daemon now takes optional extra-args; register-daemon!, maybe-register-daemon! pass them through. - *secp-daemon-extra-args* = '("--window-w" "4") activates v3 in the spawned secp256k1 daemon. cuda-shake-fanout and cuda-bignum-cgbn spawn unchanged. Deployed: - 3090-ai: rebuilt secp256k1-batch-mul, restarted via /tmp/launch.sh. Worker log shows [v3-window-w4] on smoke at n=16 and n=200. - ai.foxhop.net (4090): same. sm_89, libgmp at ~/local/gmp. Smoke test: small BSCP request to each worker over TCP, byte-identical vs coincurve. Both PASS. Catalog (CATALOG.md, www/bend.html): promoted v3 throughput to 13.83 Mkeys/s @ n=1M, recorded v4 regression and warp-scan Phase B/D as Day-5+ refactor. Progress doc: plans/form-A-day4-progress.md. --- examples/cuda-fanout/CATALOG.md | 2 +- examples/cuda-fanout/gpu-worker.lsp | 41 ++++++-- .../cuda-fanout/plans/form-A-day4-progress.md | 98 +++++++++++++++++++ examples/cuda-fanout/secp256k1-batch-mul.cu | 96 +++++++++++++++++- www/bend.html | 2 +- 5 files changed, 223 insertions(+), 16 deletions(-) create mode 100644 examples/cuda-fanout/plans/form-A-day4-progress.md diff --git a/examples/cuda-fanout/CATALOG.md b/examples/cuda-fanout/CATALOG.md index 4e096be..00bef02 100644 --- a/examples/cuda-fanout/CATALOG.md +++ b/examples/cuda-fanout/CATALOG.md @@ -22,7 +22,7 @@ form-status column says `planned` until numbers exist. | `cuda-shake-fanout` | `shake256-fanout` | 3090 | 12x over host hashlib at 1M x 16 B inputs | `(cuda-shake-fanout (quote (hex …)) out-bytes)` + BSHK binary | | `cuda-sim-ops-bin` | `ecdsa/cuda/demo_ops` | 3090 | 1.07x at 128 batches (8192 shots); crossover ~115 batches | `(cuda-sim-ops-bin "path/to/ops.bin" n-batches)` | | `cuda-bignum-cgbn` | `cgbn-batch-worker` | 3090 | 1.28 Gops/s kernel mod-mul at n=1M (256-bit, ~256x GMP CPU); 9 ops total | `BCGB` binary: op_id + bitwidth + n + modulus + a + b | -| `cuda-secp256k1-batched-mul`| `secp256k1-batch-mul` | 3090 | v1 7.87 Mkeys/s @ n=1M; v3 (windowed-G w=4) **13.60 Mkeys/s @ n=1M** (1.73x v1; ~309x coincurve); Day-4 v3 + Montgomery batch inv stacking pending | `BSCP` binary: scalars + base-point | +| `cuda-secp256k1-batched-mul`| `secp256k1-batch-mul` | 3090 | v1 7.86 Mkeys/s @ n=1M; v3 (windowed-G w=4) **13.83 Mkeys/s @ n=1M** (1.76x v1; ~309x coincurve); Day-4 v4 (v3 ladder + Montgomery batch inv) regressed -12% vs v3 due to v2 serial-per-block walks; warp-scan Phase B/D refactor needed before v4 wins; daemon default flipped to `--window-w 4` | `BSCP` binary: scalars + base-point | | `cuda-sim-axis-flip` | `ecdsa/cuda/demo_axis` | 3090 | 217 Mops/s @ K=32 M=4 (per-candidate parallel); 23.7x over per-shot N=4 at same M; LOSES to per-shot by 14% at full N=128 saturation | `(cuda-sim-axis (variant-paths …) n-shots)` | `cuda-sim-ops-bin` numbers are humbling on purpose: kickmix is diff --git a/examples/cuda-fanout/gpu-worker.lsp b/examples/cuda-fanout/gpu-worker.lsp index a81e7a6..218b2ff 100644 --- a/examples/cuda-fanout/gpu-worker.lsp +++ b/examples/cuda-fanout/gpu-worker.lsp @@ -72,8 +72,14 @@ ;;; needed is here. Should land in each tier's primitive table ;;; alongside `tcp-*`. -(define (start-daemon binary-path) - (let* ((pair (spawn-process-stdio binary-path '("--daemon"))) +;; start-daemon optionally takes extra-args (a list of strings) appended +;; after "--daemon" in the spawn argv. Day-4: secp256k1 daemon flips to +;; pass '("--window-w" "4") because v3 (Day-3 windowed-G ladder) clean- +;; wins at n>=100k vs v1; see plans/form-A-day4-progress.md. +(define (start-daemon binary-path . extra-args) + (let* ((argv (cons "--daemon" + (if (null? extra-args) '() (car extra-args)))) + (pair (spawn-process-stdio binary-path argv)) (ready (read-line (cdr pair)))) (if (string=? ready "ready") pair @@ -99,11 +105,18 @@ (else (cons 'error (string-append "?: " line))))) (else (cons 'error (string-append "?: " line))))))) -(define (register-daemon! op-name binary-path) - (set! *daemons* - (cons (cons op-name (start-daemon binary-path)) *daemons*)) +(define (register-daemon! op-name binary-path . extra-args) + (let ((pair (if (null? extra-args) + (start-daemon binary-path) + (start-daemon binary-path (car extra-args))))) + (set! *daemons* + (cons (cons op-name pair) *daemons*))) (display "gpu-worker: ready ") (display op-name) - (display " <- ") (display binary-path) (newline)) + (display " <- ") (display binary-path) + (cond + ((not (null? extra-args)) + (display " extra-args=") (display (car extra-args)))) + (newline)) ;;; -- op handler -- cuda-shake-fanout ---------------------------- @@ -382,21 +395,31 @@ ;; Optional registration — only spawn the daemon when its binary is ;; reachable. Lets a worker host serve a subset of forms without ;; failing to start because some bend form's daemon isn't installed. -(define (maybe-register-daemon! op-name binary-path) +(define (maybe-register-daemon! op-name binary-path . extra-args) (cond ((file-exists? binary-path) - (register-daemon! op-name binary-path)) + (if (null? extra-args) + (register-daemon! op-name binary-path) + (register-daemon! op-name binary-path (car extra-args)))) (else (display ";;; gpu-worker: skipping ") (display op-name) (display " - binary not found at ") (display binary-path) (newline)))) +;; Day-4 daemon flag set for the secp256k1 binary. v3 (--window-w 4) +;; landed 13.83 Mkeys/s @ n=1M (1.73x v1, 1.55x v4). Day-4 v4 stack +;; regressed -12% vs v3; daemon stays on v3 until warp-scan Phase B/D +;; lands (see plans/form-A-day4-progress.md). +(define *secp-daemon-extra-args* '("--window-w" "4")) + (define (main) (let ((port (parse-port-arg *argv*))) (set! *worker-port* port) (maybe-register-daemon! 'cuda-shake-fanout *binary-shake-fanout*) (maybe-register-daemon! 'cuda-bignum-cgbn *binary-cgbn-batch*) - (maybe-register-daemon! 'cuda-secp256k1-batched-mul *binary-secp256k1-batch*) + (maybe-register-daemon! 'cuda-secp256k1-batched-mul + *binary-secp256k1-batch* + *secp-daemon-extra-args*) (let ((server (tcp-listen port))) (cond ((eq? server #f) diff --git a/examples/cuda-fanout/plans/form-A-day4-progress.md b/examples/cuda-fanout/plans/form-A-day4-progress.md new file mode 100644 index 0000000..ceffe77 --- /dev/null +++ b/examples/cuda-fanout/plans/form-A-day4-progress.md @@ -0,0 +1,98 @@ +# Form A Day-4 progress — v3+v2 stack lands byte-identical but regresses 12% at n=1M; daemon flips to v3 + +Status: **v4 kernel landed correctness-clean, throughput-regressed. Daemon default flips to v3 (`--window-w 4`), v3 stays canonical worker path.** v3 remains best-measured kernel; v4 stays in tree as future warp-scan launchpad. + +## What shipped + +`secp256k1-batch-mul.cu` now carries four kernel variants: + +- `secp_mul_batch_kernel_v1` (Day-1) — per-thread Jacobian double-and-add, per-thread `_ModInv`. +- `v2_phase_a..d` (Day-2 batch-inv pipeline, binary ladder Phase A). +- `secp_mul_batch_kernel_v3` (Day-3) — per-thread windowed-base ladder, per-thread `_ModInv`. +- `v4_phase_a` (Day-4) — windowed-base ladder Phase A, **stacked with v2_phase_b/c/d for batch inversion**. + +CLI matrix: + +| flags | variant | +|-----------------------------------|---------| +| (none) | v1 | +| `--batch-inv` | v2 | +| `--window-w 4` | v3 | +| `--window-w 4 --batch-inv` | v4 | + +## Validation — byte-identical at every N + +Coincurve oracle on 3090-ai, default seed 0xC0FFEE: + +``` +v4 known-small n=8 PASS [v4-window-w4+batch-inv] +v4 random n=32 PASS +v4 random n=1000 PASS +v4 random n=10000 PASS +v4 random n=100000 PASS +v4 random n=1000000 PASS +``` + +Math is correct across all four variants. The HARD RULE on byte-identity holds; v4 commits because the kernel ships clean. + +## Throughput — v4 regresses 12% vs v3 at n=1M + +3090-ai best-of-3, kernel ms (after one warm-up): + +| n | v1 ms | v1 Mkeys/s | v3 ms | v3 Mkeys/s | v4 ms | v4 Mkeys/s | v4/v1 | v4/v3 | +|-----------|--------|------------|--------|------------|--------|------------|-------|-------| +| 10,000 | 2.53 | 3.95 | 4.64 | 2.16 | 5.59 | 1.79 | 0.45x | 0.83x | +| 100,000 | 13.45 | 7.44 | 10.84 | 9.22 | 12.42 | 8.05 | 1.08x | 0.87x | +| 1,000,000 | 127.29 | 7.86 | 72.32 | 13.83 | 82.27 | 12.16 | 1.55x | 0.88x | + +v4 beats v1 at n >= 100k (1.55x at n=1M) but **regresses 12% vs v3 at n=1M**. HARD-RULE threshold was "v4 must beat v3 by >= 10%". v4 fails that threshold → no v4 promotion to daemon default. + +## Why v4 lost to v3 + +Day-3 progress doc predicted: "Once v3 lands, the bottleneck shifts BACK toward scalar_mul's residual cost + per-thread inversion. The Day-2 batch-inv idea now becomes viable as a v3+v2 stacked variant." That prediction was half-right. The stacked v4 does shave the per-thread ModInv cost (~9 ms of work disappears at n=1M, because Phase B + C + D collectively run in ~10 ms vs v3's ~73 ms total). But v2's Phase B/D walks add ~20 ms of single-thread-per-block serial work that exceeds the inversion savings. + +Concretely at n=1M with v2_block_size=256: +- Phase A (v4 windowed scalar_mul, all 1M threads parallel): ~62 ms — same as v3's ladder cost. +- Phase B (3,906 blocks, 1 thread each, per-block 256-mult serial): ~7-8 ms. +- Phase C (1 thread total, 3906-block walk + 1 ModInv): ~1 ms. +- Phase D (3,906 blocks, 1 thread each, per-block 256-mult + 256-affine-proj serial): ~12 ms. + +Total ~82 ms vs v3's 72 ms. The single ModInv saved is roughly 0.3 ms (one ModInv on a 3090 takes microseconds when not blocking other threads). The per-block serial walks dwarf that saving. + +This is exactly the failure mode Day-2 progress doc called out: "Phase B and Phase D use 1 thread per block. At n=1M that's 3906 active threads vs v1's 1M threads. 3090 has 82 SMs x 1536 max-resident-threads = ~125k concurrent threads; v2 leaves >97% of SM throughput idle in the prefix and suffix passes." + +v3's per-thread ModInv saturates 1M parallel threads on those 82 SMs. The math wins. The threading model loses. + +## What v4 needs to actually beat v3 + +The same Day-2 follow-up that was deferred then: + +**Warp-level prefix scan inside Phase B + suffix walk in Phase D.** Kogge-Stone or Sklansky scan with 32 threads/warp. At n=1M with block_size=256, each warp would scan 8 elements — turn the serial 256-mult chain into 8 mults of warp-coordinated work. Estimated Phase B cost drops from ~7 ms to <1 ms; Phase D drops similarly. That should bring v4 to ~65 ms total, beating v3's 72 ms. + +That refactor is multi-day work (cross-warp shuffles, careful boundary handling, lane-mask discipline). Defer to Day-5 if fox keeps the v4 frontier alive. + +## Daemon flip — v3 wins, daemon default now `--window-w 4` + +v3 vs v1 at n=1M = 1.73x. That meets the daemon-flip threshold. Daemon spawns now carry `--window-w 4`. + +Two options were sketched in task brief: +- **Option A**: extend `start-daemon` arglist (minimal code change in lumbda). +- **Option B**: wrap binary in a `/tmp/secp-daemon-wrapper.sh` (per-host configurable). + +Chose **Option A**: add a `secp-extra-args` parameter to `gpu-worker.lsp`, modify `maybe-register-daemon!` for the secp daemon to pass `'("--daemon" "--window-w" "4")`. Keeps all configuration inside the lumbda source — single source of truth. + +## Files touched + +- `examples/cuda-fanout/secp256k1-batch-mul.cu` — added `v4_phase_a` kernel template + v4 dispatch branch in `process_one_bin` + variant label + usage string. +- `examples/cuda-fanout/gpu-worker.lsp` — added `start-daemon` arglist override, plumbed `--window-w 4` into the secp256k1 daemon spawn. +- `examples/cuda-fanout/CATALOG.md` — promoted form A throughput number from v3 13.60 Mkeys/s to v3-confirmed 13.83 Mkeys/s @ n=1M. +- `www/bend.html` — same throughput promotion. + +## Rebuild + deploy + +``` +ssh 3090-ai.foxhop.net 'cd /home/fox/git/lumbda/examples/cuda-fanout && rm -f secp256k1-batch-mul && make secp256k1-batch-mul' +ssh ai.foxhop.net 'cd /home/fox/git/lumbda/examples/cuda-fanout && rm -f secp256k1-batch-mul && make secp256k1-batch-mul' +rsync gpu-worker.lsp -> both +bash /tmp/launch.sh on both +``` diff --git a/examples/cuda-fanout/secp256k1-batch-mul.cu b/examples/cuda-fanout/secp256k1-batch-mul.cu index 45a98e8..a726d21 100644 --- a/examples/cuda-fanout/secp256k1-batch-mul.cu +++ b/examples/cuda-fanout/secp256k1-batch-mul.cu @@ -590,6 +590,63 @@ __global__ void secp_mul_batch_kernel_v3( limbs_to_be32(outp + 32, Ay); } +/* ------------------------------------------------------------------ */ +/* Kernel v4: Day-4 stack — v3 windowed-G ladder feeds the v2 */ +/* Montgomery batch-inversion pipeline. */ +/* */ +/* Pipeline: */ +/* Phase A (v4_phase_a): per-thread windowed scalar_mul, writes */ +/* Jacobian (Jx, Jy, Jz) to scratch[idx*12..idx*12+12) and writes */ +/* z[idx] = Jz (or 1 if Jz == 0) to scratch z[]. */ +/* Phase B (v2_phase_b reused): per-block prefix product over z[], */ +/* writes c[] and block_prod[]. */ +/* Phase C (v2_phase_c reused): single-thread cross-block walk, */ +/* ONE _ModInv on the full product, derives G_left[] + u_entry[]. */ +/* Phase D (v2_phase_d reused): per-block backward walk, emits */ +/* affine BE bytes. */ +/* */ +/* The v3 win (cut scalar_mul cost by ~16% of field-mults) compounds */ +/* with the v2 win (collapse N parallel ModInvs to 1 + N-1 ModMults) */ +/* because Day-3 progress doc identified: "Once v3 lands, the */ +/* bottleneck shifts BACK toward scalar_mul's residual cost + */ +/* per-thread inversion." v4 retires both bottlenecks in one pipeline. */ + +template +__global__ void v4_phase_a( + const uint64_t* base_xy, /* unused but kept for ABI parity */ + const uint64_t* table_g, /* W * 8 u64 in global memory */ + const uint8_t* scalars, + uint64_t* jac, /* n * 12 u64: [Jx(4), Jy(4), Jz(4)] */ + uint64_t* z, /* n * 4 u64: Jz, or 1 if Jz == 0 */ + uint32_t n, + uint32_t w) +{ + __shared__ uint64_t s_table[W * 8]; + for (uint32_t i = threadIdx.x; i < W * 8; i += blockDim.x) { + s_table[i] = table_g[i]; + } + __syncthreads(); + + uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx >= n) return; + const uint8_t* k_be = scalars + (size_t)idx * 32; + + uint64_t Jx[4], Jy[4], Jz[4]; + scalar_mul_windowed(Jx, Jy, Jz, k_be, s_table, w, W); + + uint64_t* slot = jac + (size_t)idx * 12; + slot[0]=Jx[0]; slot[1]=Jx[1]; slot[2]=Jx[2]; slot[3]=Jx[3]; + slot[4]=Jy[0]; slot[5]=Jy[1]; slot[6]=Jy[2]; slot[7]=Jy[3]; + slot[8]=Jz[0]; slot[9]=Jz[1]; slot[10]=Jz[2]; slot[11]=Jz[3]; + + uint64_t* zslot = z + (size_t)idx * 4; + if (is_zero4(Jz)) { + zslot[0] = 1; zslot[1] = 0; zslot[2] = 0; zslot[3] = 0; + } else { + zslot[0] = Jz[0]; zslot[1] = Jz[1]; zslot[2] = Jz[2]; zslot[3] = Jz[3]; + } +} + /* ------------------------------------------------------------------ */ /* Kernel v2: Montgomery batch inversion across N independent z's. */ /* */ @@ -1122,7 +1179,32 @@ static int process_one_bin(const char* in_path, const char* out_path) { struct timespec t0, t1; clock_gettime(CLOCK_MONOTONIC, &t0); - if (g_window_w > 0) { + if (g_window_w > 0 && g_use_batch_inv) { + /* v4: v3 windowed scalar_mul (Phase A) + v2 batch inversion */ + /* (Phase B/C/D). Stacked Day-4 path. */ + CUDA_CHECK(cudaMalloc(&d_table, (size_t)v3_W * 8 * sizeof(uint64_t))); + CUDA_CHECK(cudaMalloc(&d_jac, (size_t)n * 12 * sizeof(uint64_t))); + CUDA_CHECK(cudaMalloc(&d_z, (size_t)n * 4 * sizeof(uint64_t))); + CUDA_CHECK(cudaMalloc(&d_c, (size_t)n * 4 * sizeof(uint64_t))); + CUDA_CHECK(cudaMalloc(&d_block_prod, (size_t)v2_K * 4 * sizeof(uint64_t))); + CUDA_CHECK(cudaMalloc(&d_G_left, (size_t)v2_K * 4 * sizeof(uint64_t))); + CUDA_CHECK(cudaMalloc(&d_u_entry, (size_t)v2_K * 4 * sizeof(uint64_t))); + + v3_table_init<<<1, 1>>>(d_base, d_table, v3_W); + if (g_window_w == 4) { + v4_phase_a<16><<>>( + d_base, d_table, d_scalars, d_jac, d_z, n, 4u); + } else { + uint32_t tpb8 = 64; + uint32_t blocks8 = (n + tpb8 - 1) / tpb8; + v4_phase_a<256><<>>( + d_base, d_table, d_scalars, d_jac, d_z, n, 8u); + } + v2_phase_b<<>>(d_z, d_c, d_block_prod, n, v2_block_size); + v2_phase_c<<<1, 1>>>(d_block_prod, d_G_left, d_u_entry, v2_K); + v2_phase_d<<>>(d_jac, d_z, d_c, d_G_left, d_u_entry, + d_points, n, v2_block_size); + } else if (g_window_w > 0) { /* v3: precompute (i*P) table, then run windowed-ladder kernel. */ CUDA_CHECK(cudaMalloc(&d_table, (size_t)v3_W * 8 * sizeof(uint64_t))); v3_table_init<<<1, 1>>>(d_base, d_table, v3_W); @@ -1190,10 +1272,13 @@ static int process_one_bin(const char* in_path, const char* out_path) { int wr = write_buffer_to_file(out_path, out_buf, out_size); free(out_buf); const char* variant_label; - char v3_label[32]; - if (g_window_w > 0) { - snprintf(v3_label, sizeof(v3_label), "v3-window-w%u", g_window_w); - variant_label = v3_label; + char vN_label[40]; + if (g_window_w > 0 && g_use_batch_inv) { + snprintf(vN_label, sizeof(vN_label), "v4-window-w%u+batch-inv", g_window_w); + variant_label = vN_label; + } else if (g_window_w > 0) { + snprintf(vN_label, sizeof(vN_label), "v3-window-w%u", g_window_w); + variant_label = vN_label; } else if (g_use_batch_inv) { variant_label = "v2-batch-inv"; } else { @@ -1295,6 +1380,7 @@ int main(int argc, char** argv) { "\n" " --window-w 4|8 Day-3 v3: windowed-base ladder (precomputed table of i*P)\n" " --batch-inv Day-2 v2: Montgomery batch inversion across N points\n" + " --window-w 4 --batch-inv Day-4 v4: v3 ladder + v2 batch inversion stacked\n" " --no-batch-inv Day-1 v1: per-thread _ModInv (default)\n" " --no-window disables window ladder (default)\n", argv[0], argv[0]); diff --git a/www/bend.html b/www/bend.html index b217727..c29fbea 100644 --- a/www/bend.html +++ b/www/bend.html @@ -143,7 +143,7 @@ make gpu-worker LUMBDA=asm # smallest footprint cuda-secp256k1-batched-mul live RTX 3090 - v1 7.87 Mkeys/s @ n=1M; v3 (windowed-G w=4) 13.60 Mkeys/s @ n=1M — 1.73× v1, ~309× coincurve CPU. Day-4 v3 + Montgomery batch inv stacking is the next step toward FixedPaul's 6.5 Gkeys/s on 4090 + v1 7.86 Mkeys/s @ n=1M; v3 (windowed-G w=4) 13.83 Mkeys/s @ n=1M — 1.76× v1, ~309× coincurve CPU. Day-4 v4 (v3 ladder + Montgomery batch inv) regressed -12% vs v3 due to v2's serial-per-block Phase B/D walks; warp-scan Phase B/D refactor needed before v4 wins. Daemon default now ships --window-w 4 BSCP binary: scalars + base-point → BSCR points