# Form A Day-3 progress — windowed-G ladder cuts scalar_mul, byte-identity holds at 1M, 1.73x at n=1M

Status: **landed on neoblanka working tree; commit pending fox approval.** v3 kernel passes byte-identical against coincurve at n in {32, 1000, 10000, 100000, 1000000} and outperforms v1 at every n >= 100k. 3090-ai bench result locked.

## What shipped

`secp256k1-batch-mul.cu` now carries three 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) — kept for future warp-scan rewrite; gated by `--batch-inv`.
- `secp_mul_batch_kernel_v3<W>` (Day-3) — per-thread windowed-base ladder, per-thread `_ModInv`. Templated on table size W = 2^window_w.

Companion device kernel `v3_table_init` builds W affine (i*P) entries from a base point in one launch (1 thread, ~15 jac-add-affine + 1 ModInv + 14 ModMult, all amortised across N).

## CLI

- `--window-w 4` selects v3 with a 16-entry table (Day-3 canonical).
- `--window-w 8` reserved but stub-rejected: device-side `v3_table_init` caps at W <= 16 to keep register/stack bounded. w=8 needs a host-provided scratch buffer; deferred.
- `--no-window` (default) keeps v1/v2 binary ladder.
- `--no-batch-inv` (now default ON) selects v1 over v2. Day-2's regression made v2 a footgun for daemon callers; v3 is the new winning path.

## Algorithm

Per scalar k:

1. Q = O.
2. For widx = 0..63 (64 windows of w=4 bits, MSB to LSB):
   - Q = 16 * Q  (four jac_doubles).
   - wv = bits 252-4*widx .. 255-4*widx of k.
   - If wv != 0: Q = Q + table_shared[wv]  (one jac_add_affine).
3. Output Q.

Cost per scalar: 256 doubles + ~63 adds (vs ~128 in binary double-and-add). ~16% fewer field-mult ops, but the win compounds because adds carry more ModMult per call than doubles do.

Table loaded cooperatively into `__shared__ uint64_t s_table[16*8]` (1024 B) at block start; per-thread divergent indexing then hits shared memory, not constant memory (which broadcasts only on uniform lookup).

## Validation — byte-identical at every N

Coincurve oracle on 3090-ai, default seed 0xC0FFEE:

```
known-small  n=8        PASS  [v3-w4]
random n=32             PASS  [v3-w4]
random n=1000           PASS  [v3-w4]
random n=10000          PASS  [v3-w4]
random n=100000         PASS  [v3-w4]
random n=1000000        PASS  [v3-w4]
```

Default v1 path (no flags) still PASSes as before — Day-3 is purely additive.

## Throughput — v3 wins at n >= 100k

3090-ai best-of-3 with `--no-batch-inv` (excluding v2 from the comparison):

| n         | v1 ms  | v1 Mkeys/s | v3 ms  | v3 Mkeys/s | v3 / v1 |
|-----------|--------|------------|--------|------------|---------|
| 10,000    | 2.37   | 4.22       | 4.58   | 2.18       | 0.52x   |
| 100,000   | 13.36  | 7.48       | 10.77  | 9.29       | 1.24x   |
| 1,000,000 | 127.04 | 7.87       | 73.54  | 13.60      | 1.73x   |

At n=10k, v3 loses because table-init kernel overhead and the larger per-thread state amortise badly across only ~10k workers. At n=100k and above, v3 dominates. At n=1M the ratio matches the theoretical adds-cut estimate plus a constant-launch-cost amortisation.

## Why v3 wins where v2 lost

Day-2 v2 traded N parallel ModInvs for one ModInv + N-stage serial Phase B/D walks. At n=1M, Phase B/D ran 1 thread per block over 3906 blocks (vs v1's 1M concurrent threads) — total kernel time dropped a little for the inversion portion but the entire kernel grew by the per-block serial multiplies. v3 leaves all N threads fully parallel; the only serial work is the constant-cost 1-thread table init, which is sub-millisecond.

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** (Day-4 candidate): build the windowed table once, run v3's per-thread loop, then collect all N Z values and feed them into v2's Montgomery suffix walk. That stacking is the proper terminus per plan section 10.

## Files touched

- `examples/cuda-fanout/secp256k1-batch-mul.cu` — added `v3_table_init`, `scalar_mul_windowed`, `secp_mul_batch_kernel_v3<W>`, host wiring for `--window-w` flag, variant label in stderr. Default `g_use_batch_inv` flipped from 1 to 0 (v2 was Day-2 default but regressed; Day-3's v3 is the new canonical path under `--window-w 4`; default no-flag behaviour stays v1).

## Rebuild path

```
ssh 3090-ai.foxhop.net 'cd /home/fox/git/lumbda/examples/cuda-fanout && rm -f secp256k1-batch-mul && make secp256k1-batch-mul'
./secp256k1-batch-mul --window-w 4 --binary in.bscp out.bscr
```

## Open question for fox

Should the gpu-worker.lsp daemon spawn pass `--window-w 4` by default? Right now the daemon receives no flags, so it runs v1. To make v3 the canonical worker path, we'd add `--window-w 4` to the spawn argv in `gpu-worker.lsp`. Holding that change until fox confirms the bench numbers reproduce in his hands.
