engineering/lumbda
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

examples/cuda-fanout: Day-4 v4 stack landed (byte-clean, -12% vs v3); daemon flips to v3 --window-w 4

Browse source 
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.
This commit is contained in:
russell@unturf.com 2026-06-05 19:07:00 -04:00
parent 3e0ca555b5
commit 3a8c66a148
No known key found for this signature in database
5 changed files with 223 additions and 16 deletions
Download patch file Download diff file Expand all files Collapse all files

2
examples/cuda-fanout/CATALOG.md
View file

@ -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

41
examples/cuda-fanout/gpu-worker.lsp
View file

@ -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)

98
examples/cuda-fanout/plans/form-A-day4-progress.md Normal file
View file

@ -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<W>` (Day-3) — per-thread windowed-base ladder, per-thread `_ModInv`.
- `v4_phase_a<W>` (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<W>` 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
```

96
examples/cuda-fanout/secp256k1-batch-mul.cu
View file

@ -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<W>): 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 <uint32_t W>
__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><<<blocks, threads_per_block>>>(
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><<<blocks8, tpb8>>>(
d_base, d_table, d_scalars, d_jac, d_z, n, 8u);
}
v2_phase_b<<<v2_K, 1>>>(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<<<v2_K, 1>>>(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]);

2
www/bend.html
View file

@ -143,7 +143,7 @@ make gpu-worker LUMBDA=asm # smallest footprint</code></pre>
<td><code>cuda-secp256k1-batched-mul</code></td>
<td>live</td>
<td>RTX 3090</td>
<td>v1 7.87 Mkeys/s @ n=1M; <strong>v3 (windowed-G w=4) 13.60 Mkeys/s @ n=1M</strong> &mdash; 1.73&times; v1, ~309&times; coincurve CPU. Day-4 v3 + Montgomery batch inv stacking is the next step toward FixedPaul's 6.5 Gkeys/s on 4090</td>
<td>v1 7.86 Mkeys/s @ n=1M; <strong>v3 (windowed-G w=4) 13.83 Mkeys/s @ n=1M</strong> &mdash; 1.76&times; v1, ~309&times; 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 <code>--window-w 4</code></td>
<td><code>BSCP</code> binary: scalars + base-point &rarr; <code>BSCR</code> points</td>
</tr>
<tr>