Each gpu-worker.lsp now listens on both wire-TCP (existing :8320) and
HTTP/1.1+CORS (new :8321), sharing one handle-request dispatcher. Lets
a tab on https://lumbda.com/playground/ POST to its own machine via
http://localhost:8321/ — browsers permit localhost from HTTPS origins
without TLS, so no proxy, no cert, no fox-owned infra required for the
decentralized run-your-own-bend story.
main() forks at startup: child runs http-run-loop on :8321, parent
keeps existing run-loop on :8320. Adding a new op-head to handle-request
exposes it over both transports automatically. Binary modes
(BSHK/BCGB/BSCP/BSRT/BSB3) stay wire-only — they exist for native
callers who already cache the binary locally; browser callers send
S-expression recipes the worker dispatches the same way.
Two latent defects fixed to make CPU-only and Python-tier hosts work:
- vram-used-mib now file-exists? guards /usr/bin/nvidia-smi. Python
tier's spawn-process-stdio raises FileNotFoundError on missing
binary, not returning #f as the prior code expected, which crashed
every worker on a CPU-only laptop.
- fork-self return discriminated via (number? pid) not (eq? pid 0).
Python tier's (eq? 0 #f) returns #t because == conflates int 0
with bool False; pre-existing run-loop has the same risk but
C/asm tier (identity eq?) masks it for the production case.
Phase 2 (server-side factory ops: compile uploaded .lsp recipes into
.bin before bending — the foxhop champion-circuit workflow) deferred
until authentication lands; today a worker on the public internet
would let any caller occupy our GPU.
Operational Caddy + DNS proposals in plans/bend-http-deploy.md cover
the personal-remote-access endpoint chain (proxy.unturf.com edge →
ai.foxhop.net Caddy → 3090-ai:8321) gated by trusted-IP allowlist —
applied separately.
Also codifies the playground "CSS Grid only, never flexbox" rule in
CLAUDE.md: all www/ and wasm/ stylesheets are already grid-only;
documenting the invariant so future edits don't drift.
Tests: smoke-bend-http.sh — (ping)→(ok pong), unknown-op fallback,
OPTIONS CORS preflight — all PASS. Wire path unchanged, verified
round-trip via 8-digit-prefix framing.
Two link fixes:
- Drop dead /gpu-mesh link. That wiki RST exists in our private
repo but our remarkbox host has not landed it under a public URL
yet. Describe coordinator pattern inline so a reader sees the
shape without following an unreachable href.
- /ecdsa link was also wrong. Real published page lives at our
remarkbox UUID slug. Updated to actual URL.
Closed-source reminder for future drift: foxhop research code is
NOT open source. Never link to a script path under foxhop-book; if a
concept needs explaining, inline a description instead of pointing
at our private repo.
Earlier text linked to 'foxhop ecdsa scripts/coordinator-mesh.py' as
if reachable at a public URL. That script lives only in the private
foxhop-book repo — no public path exists. Point readers at our
foxhop.net/gpu-mesh wiki page instead (lives in remarkbox, publicly
reachable, describes the cost-routed coordinator pattern at concept
level).
Previous override example dropped to PORT=9001 — drifted off our BEND
mnemonic. Updated so LUMBDA=python & LUMBDA=asm overrides keep our
default port. Off-BEND port override stays available for a 2nd worker
on a single host, called out as such.
Extends the cuda-sim-ops-bin row in our form catalog with a measured
property: kernel wall scales linearly with Σ Toffoli across circuit
variants.
foxhop ecdsa lever sweep at p=251 (n+1=9 secp256k1-toy) dispatched 6
lumbda-emitted ops.bin variants through demo_ops on a 3090. Wall vs
Toffoli per shot:
Fermat-schoolbook 167984 Tof/shot 261.4 ms/batch
refined-Solinas 24176 Tof/shot 26.6 ms/batch
9.83x GPU wall reduction matches 6.95x Toffoli ratio + parallel
Clifford drop. No scheduler surprise — pick any QECCOPS1 circuit,
predict GPU wall from a cheap CPU op-counter.
Links to https://www.foxhop.net/ecdsa for the upstream lever-attack
context.
demo_ops in foxhop ecdsa repo now defaults --rng-mode shake (Fiat-Shamir
over op stream, ports eval_circuit::fiat_shamir_seed verbatim). Σ Toffoli
& avg Toffoli on stock ops.bin match upstream eval_circuit bit-for-bit:
15,999,651,264 / 1,773,011.000 across 141 batches (9024 shots) on a
3090. Our cuda-sim-ops-bin numbers now compare directly to upstream's
public Pareto scoreboard with no calibration constant.
Catalog row & bend.html row updated to document the shake/lfsr toggle
& the upstream-match guarantee. Throughput figure also refreshed to
1.27x at 141 batches (from prior 1.07x at 128 batches).
Documents the (health) op + *worker-health* cache + VRAM-ranked
pick shipped in lumbda 58fd787. Three subsections:
- Worker side: (health) op shape with measured numbers
- Client side: cache TTL (5 s ok / 30 s down) + ranking logic
- Failure handling: dispatch errors flip workers down;
with-exception-handler wraps probes so a dead peer never
aborts a multi-worker iteration
Integrated into bend.html's fleet narrative — backward-compat
note so older workers without (health) still register as
available with vram=0.
Bullets stay technical & terse; no commit hashes, no internal
implementation quirks (eq? 0 #f / Python exception propagation
stay in our commit log).
Agent ac6c3c7b built blake3-fanout.cu (clean-room BLAKE3 reference
adaptation, vendored under BLAKE3 team's CC0-1.0 / Apache-2.0
allowance) + test_blake3_known_answers.py harness + Makefile target
+ gpu-worker.lsp handler (handle-binary-blake3, BSB3/BSR3 magic).
Bench on 3090-ai (best-of-3, kernel-only):
workload kernel throughput
1k × 64 B 0.08 ms ~7 GB/s
100k × 64 B 1.87 ms 59.7 GB/s
1M × 64 B 1.97 ms 32.5 GB/s
1k × 1 MB 51.5 ms 20.4 GB/s
Byte-identity vs the BLAKE3 reference spec PASS at n in
{32, 1k, 10k, 100k, 1M}; covers both the single-chunk (≤1024 B)
& multi-chunk (≥1024 B) Merkle-tree paths.
Hardware policy compliance: per the 3090-only directive (commit
9e4e9b4), this row lists RTX 3090 only. The agent's 4090
measurements were dropped from the catalog & bend.html. ai
worker stays disabled.
Note the agent hit a socket error AFTER all files landed but
BEFORE it could git commit. This commit assembles its work from
the working tree, verifies the build still passes on 3090
(byte-identity + bench above), and ships.
Live forms now: 7 — shake, sim-ops-bin, sim-axis-flip,
cgbn (9 ops), secp256k1 (v3 windowed-G), radix-sort,
blake3-tree.
Fox decision: don't routine-fan-out to ai.foxhop.net (4090) when
qwen LLM holds GPU residency. The radix-sort 4090 OOM caveat
surfaced today demonstrated the cost of casual co-residency —
secp daemon parked 24 GiB up front leaving 47 MiB free.
bend.lsp's *bend-workers* default already empty (single-host
fallback). Updated the docstring example to drop the ai.foxhop.net
entry; multi-host fan-out is OPT-IN per call via bend-set-workers!
or BEND_WORKERS env.
bend.html fleet section now reflects the policy:
- 3090-ai.foxhop.net:9091 active production worker
- ai.foxhop.net:9092 reserved for qwen; bend per workload
ai.foxhop.net worker process killed; 4090 VRAM returned to qwen
(1.6 GiB free post-kill vs 47 MiB while bend was running).
When we have a long-running parallel sweep that justifies the
4090's marginal throughput, the caller opts in explicitly. Don't
auto-route.
Thin CUDA binary wrapping cub::DeviceRadixSort::SortKeys on a 64-bit
key stream. One op wired day-1 (0x01 sort-u64-asc); 0x02/0x03/0x04
slots reserved (desc, u32, key-value) for future builds.
Wire stays distinct from existing forms:
request: BSRT | u32 op_id | u32 n | u64[n]
response: BSRR | u32 status | u32 n | u64[n] sorted asc
Validated on 3090-ai.foxhop.net byte-identical to Python sorted() at
n ∈ {32, 1k, 100k, 1M, 10M}. Bench at sustained throughput:
n kernel_ms Gkeys/s
100,000 0.142 0.706
1,000,000 0.265 3.767
10,000,000 1.817 5.504
~4x over the published Titan baseline (1.4 Gkeys/s) at saturation,
matching CUB's expected Ampere scaling.
gpu-worker.lsp learns handle-binary-sort + BSRT magic dispatch +
maybe-register-daemon! for cuda-radix-sort (overridable via
RADIX_SORT_WORKER env). Both 3090-ai (:9091) & ai (:9092) workers
restarted; both log `ready cuda-radix-sort <- ./radix-sort`.
4090 (ai.foxhop.net) standalone --binary run OOMs on cudaMalloc when
all four daemons are co-resident (secp256k1 daemon parks ~24 GiB on
startup, leaving 47 MiB free). Pre-existing capacity constraint of
the ai host, not a form-G defect; tracked in form-G-progress.md.
CATALOG.md & www/bend.html live-forms table updated with measured
3090 numbers; Wave 1 surveyed row for G marked as promoted.
Wave 1's 7 detailed entries (3 paragraphs each) become one compact
4-column table matching Waves 2/3/4. Drops ~30 lines without
losing any reference URL or speedup number.
Protocol section: 2 paragraphs → 1.
Fleet section: 4 paragraphs + table → 2 paragraphs + same table.
Net change: 284 → ~244 lines while preserving every citation and
every measurement. Reads more like a reference doc.
Page had accreted session debugging (Form D structural finding,
opt 1/opt 3 narrative, Day-4 v4 regression notes, robustness fix
narrative, first measured sweep paragraphs). Stripped all of it.
What stays:
- what bend is
- start a worker
- call it
- wire protocol (S-exp vs binary modes + measured table)
- fleet (2-host cluster table + brief coexistence note + cluster
aggregate numbers)
- catalog: live forms (5) + Wave 1 (A-G detailed) + Wave 2 + 3
(compact tables) + why-a-form-earns-its-slot criteria
- source pointers
What goes (back to CATALOG.md / plan docs where they belong):
- Form D structural finding paragraph
- Form D opt 1 + opt 3 RESULTS narrative
- Day-4 v4 regression detail in the secp row
- First measured ECDSA-mission sweep paragraphs
- Robustness gap fix paragraph
Page length: 371 → ~226 lines. Reads as a reference doc again
instead of a session log.
Earlier commit incorrectly named ports 9090+9091 as qwen LLM serving.
Agent's actual finding (ai.foxhop.net second-worker deployment) was
uid-992 Erlang Cowboy services on those ports — unsandbox infra,
not qwen. Qwen runs on the GPU itself, not on those TCP ports.
The coexistence story (our bend kernels share a 4090 with qwen
because we burn SMs for ms-at-a-time then release) stays intact;
the false port attribution removed.
Adds a new "Fleet" section between the ecdsafail workload table &
the catalog. Documents what now runs in production:
- 2-host LAN cluster: 3090-ai:9091 + ai.foxhop.net:9092 (4090, sm_89)
- 4090 box shares the GPU with a qwen LLM (vllm) — our bend kernels
fit alongside because secp256k1 / CGBN workloads burn SMs only for
milliseconds at a time, then release
- Cross-GPU parity at small N: 0.07-0.10 ms kernel on both 3090 &
4090 for CGBN n=100k mod-mul; both starve between calls so the
4090's FLOPS advantage doesn't show until kernels run long enough
to amortize wire overhead
- Cluster aggregate: 1.6× speedup at 200 × n=1k (cap is client-side
serialization, not workers); async fan-out unlocks the remaining 2×
- First measured sweep through full lumbda → emit-ops-bin → bend →
demo_ops pipeline: refined Bernstein-Yang variant cuts Σ Toffoli
by 36.5% at p=11, within 3 pts of the p=251 reference prediction
(-40%). Small-fixture screen confirmed faithful predictor.
- Robustness gap fixed (handle-binary-* delete-file unconditional);
12 guard lines across 6 sites; zero crashes since deployment.
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.
foxhop CLAUDE.md update: www.foxhop.net is private (book.foxhop.net
source + research artifacts). URLs like https://www.foxhop.net/ecdsa.html
do not exist & must not appear in public posts or external links.
Public-facing copy of ECDSA work belongs on foxhop.net wiki (remarkbox
RST pages), not as static HTML under www.foxhop.net/.
Stripped two external links:
bend.html — "secp256k1 point-addition challenge work at
foxhop.net/ecdsa" → "(research artifact, internal repo)"
index.html — same pattern → "internal ECDSA research"
Also brought index.html catalog highlight in sync with Day-3:
was: secp256k1 batched point-mul (6.5 Gkeys/s on a 4090, surveyed)
now: 13.6 Mkeys/s on a 3090, ~309x coincurve CPU
The 4090 projection moves out of the lead because we now have a
measured number on our actual hardware.
Form A Day-3 shipped at lumbda ecfe27a. Windowed-G ladder (w=4,
16-entry G-table built once via batch-inverse on the 15 Z-coords)
landed clean.
Bench on 3090, n=1M (best-of-3, --no-batch-inv):
v1 (Day-1) 127.04 ms → 7.87 Mkeys/s
v3 (Day-3 w=4) 73.54 ms → 13.60 Mkeys/s (1.73x v1, ~309x coincurve)
Byte-identity PASS at n in {32, 1000, 10000, 100000, 1000000} plus
known-small edge cases (k in {1, 2, 3, 7, 0xdeadbeef, n-1, n, 2^128-1}).
Daemon default still serves v1; --window-w 4 flag selects v3
explicitly. Fox's call on flipping the daemon default.
Day-4 plan: stack v3 + Day-2 batch-inv. v3 cut scalar_mul; the
residual inversion cost now actually matters, which is what Day-2
needed to win.
Three agents still in flight: walker promotion (#39), lever
generator (#40), ai.foxhop.net second worker (#30).
Form D Option 3 (QECCOPS2 packed op format) shipped at foxhop
commit 90484ca. Numbers:
n_batches unpacked packed speedup
16 5842 ms 5428 1.076x
64 6226 ms 5830 1.068x
128 6604 ms 6201 1.065x
On-disk shrink: 716 MB → 307 MB (2.33x). Storage win, not compute win.
Critical diagnostic correction from the opt 3 agent: the 3.5x
bandwidth projection was WRONG because ops_loader.c already
narrowed u64 → u32 on load, so in-VRAM Op was already 28 B, not
56 B. Max realistic VRAM-reduction was 28 → 24 B = 1.17x best
case. The measured 1.07x matches: per-shot state traffic (qubits
+ bits per thread) is ~85x larger than the op stream, which
warp-broadcasts through L1/L2. Op-stream bandwidth was not the
bottleneck.
Both Form D pivots (opt 1 axis-flip, opt 3 packing) converged on
the same finding: the 3090 is compute-saturated at ~250 Mops/s on
the kickmix circuit, not bandwidth-saturated. Algorithm + layout
tweaks already extracted; the real next macro-lever is multi-GPU
fan-out across the fleet (3090-ai + ai.foxhop.net 4090 + future
nodes, each saturating its own ceiling in parallel).
cuda-sim-axis-flip remains LIVE as a tool with specific use:
many-candidates × few-shots search-loop early-screen. ops.bin
packing remains useful for fleet rsync (3.5x smaller payloads
across the LAN matters when shipping candidate variants).
Form D opt 1 (axis-flip refactor) shipped at foxhop 1f7ac9d.
Per-candidate parallelism over per-shot delivers 217 Mops/s at
K=32 candidates × M=4 shots on a 3090, byte-identical with the
CPU reference at every (K, M) pair we measured.
Key finding inside the numbers: the bandwidth-bound diagnosis
flagged in the Form D structural-finding doc was WRONG. Both
axes (per-shot N=128 and per-candidate K=32 M=4) saturate at the
SAME ~220-250 Mops/s on the 3090. Per-shot wins by 14% at full
saturation; axis flip wins by 23.7x at small M because it fills
SMs in one launch instead of leaving them idle.
So the axis kernel is the right tool for lumbda's many-candidates
× few-shots search-loop early-screen pattern — not a replacement
for the per-shot kernel.
Form D opt 3 (ops.bin packing) still in flight. Since compute
saturates before bandwidth on this device, packing may not
deliver expected gain. Letting that agent finish so we have
empirical numbers either way.
Live forms now: 5
cuda-shake-fanout ~12x host hashlib @ 1M
cuda-sim-ops-bin 1.07x @ 128 batches (kickmix per-shot)
cuda-bignum-cgbn 1.28 Gops/s mod-mul @ 1M (256-bit)
cuda-secp256k1-batched-mul 6.51 Mkeys/s @ 100k (148x coincurve)
cuda-sim-axis-flip 217 Mops/s @ K=32 M=4 (kickmix per-candidate)
Form A (cuda-secp256k1-batched-mul) shipped earlier this session at
commit 7661788. Form B (cuda-bignum-cgbn) extended to all 9 ops at
21bd26a. Both now serving on 3090-ai gpu-worker:9091 alongside the
shake form.
Catalog Live forms table:
cuda-bignum-cgbn | 1.28 Gops/s kernel mod-mul @ n=1M, 256-bit
(~256x GMP CPU single-thread); 9 ops total
cuda-secp256k1-batched-mul| 6.51 Mkeys/s @ n=100k (~148x coincurve CPU);
Day-2 Montgomery batch inversion projected
toward FixedPaul's 6.5 Gkeys/s on 4090
Wire magics now in use:
BSHK/BSHR — cuda-shake-fanout
BCGB/BCGR — cuda-bignum-cgbn
BSCP/BSCR — cuda-secp256k1-batched-mul
(BSTB/BSTR reserved for cuda-clifford-stabilizer; parked per
structural finding — our circuit has no H/S gates)
Form D pivots (axis-flip sim_gpu.cu + ops.bin packing) still in
flight; CATALOG.md + bend.html will gain rows once they ship.
Notable defect surfaced during Form A build: widely-cited secp256k1
generator y-coordinate
0x483ADA7726A47B0DAFFA10ED2E11458A823D0E1D89DCAB14C7C39D9F8B97C20A
does NOT satisfy y^2 = x^3 + 7 mod p. Real Gy =
0x483ADA7726A3C4655DA4FBFC0E1108A8FD17B448A68554199C47D08FFB10D4B8
Cross-tutorial propagation. UNDF candidate logged.
Form D build agent discovered: our point-add circuit contains no
Hadamard or S gates (only X/CX/CCX/CZ/CCZ/SWAP/R/HMR/Z/NEG).
State never leaves the computational basis. Aaronson-Gottesman
tableau compression buys nothing when superposition does not
exist; reduces to exactly what sim_gpu.cu already does at one
bit per qubit per shot.
Toffoli fraction measured 13.87% (well under the 40% threshold
the planner flagged). The 1.07x cuda-sim-ops-bin ceiling traces
to memory-bandwidth on per-shot striped state — not algorithm.
STABSim-class wins remain valid for QEC / surface-code work where
H + S exist; that's a future workload.
Two replacement directions queued:
1. axis-flip sim_gpu.cu — per-candidate parallelism over per-shot.
~3 days. Reuses BSHK new op_id; no AG tableau.
2. ops.bin packing — 56→16 B per op halves global-memory traffic.
Addresses the actual bottleneck.
Catalog + public bend.html both updated. Form D progress doc at
examples/cuda-fanout/plans/form-D-build-progress.md documents the
structural reasoning in full.
Splits a dense bend section out of index.html (now 6 lines: tagline,
example, three highlights, CTA) into a new public page at
lumbda.com/bend.html carrying:
* full wire protocol (S-exp + binary BSHK modes)
* tier-choice table (Py / C / asm × S-exp / binary)
* real workload table from foxhop ecdsafail measurements
* 2 live forms + 7 surveyed forms (Wave 1: secp256k1 batch-mul,
CGBN bignum, Pollard rho, Stim-on-GPU stabilizer, Bernstein-Yang
inverse, NTT, CUB sort/scan)
* 20 surveyed forms (Wave 2) sorted by speedup descending —
minhash 600-1000x, cuckoo filter 378x, ChaCha20 400 GB/s,
SAT 93x, Dilithium PQ 57.7x, BLAKE3 tree, cuFFT batched,
Bloom filter modern, GEMM FP8, hash-join 1.8T tuples/s on
1024xA100, cuGraph 38B TEPS, TRUST triangle 1T TEPS, nvCOMP 2.2x
* skipped section listing forms that don't pass the threshold,
so the catalog stays honest about what GPU dominates vs not
* cited canonical references for every entry
CATALOG.md mirrors the same data — single source of truth in repo;
bend.html renders the same metadata for public reading.
Build order remains: A cuda-secp256k1-batched-mul, D Clifford
stabilizer, B CGBN bignum.
