shopt -s nullglob (set so outer for-loop tolerates empty *.bin) made
an unmatched .inflight-* glob expand to nothing. Bare "ls >/dev/null"
then succeeded by listing CWD, and our if-branch incorrectly skipped
every orphan whose .inflight-* did not match. Net effect: medium-sized
orphan bins (100 MiB+, valid QECCOPS1 magic, no markers) accumulated
in our queue indefinitely across pool restarts.
Replace with compgen -G which returns success only when our pattern
matches, immune to nullglob.
Adds tests/integration/test-heal-orphan-bins.sh as TCRAUDT reducer
covering our contract: medium-size orphan promotes to .ready,
sub-threshold truncated to .done, bad-magic DLQs to dlq/.
Incident 2026-06-14: 8 vecC-tri-*-w8c.bin orphans (493 MiB each, Jun 12
emit) survived multiple foxhop pool restarts. Live factory queue
visible to operator as a persistent ~8-bin floor that never drained.
Adds a parallel build of all three Lumbda implementations to WASM, a
single-page playground at www/playground/, and a verified test suite.
Tiers
- Python: Pyodide (CPython-in-WASM) hosting lumbda.py
- C: Emscripten build of c/ (tree-walker + bytecode VM; jit.c
stubbed, gc.c uses its existing no-Boehm fallback)
- Asm: hand-written asm/lumbda.wat — parallel impl to asm/lumbda.s.
Reader, eval (lambda/define/if/cond/let/and/or/quote/set!),
recursion across mutated top-level env, bump allocator with
memory.grow, 24 primitives. ~1200 lines of raw WAT.
SPA (wasm/app/, deployed to www/playground/)
- CodeMirror 6 editor (Scheme highlighting) on left, output on right
- Radios: 4 demos (Mandelbrot, Fib+Ack, Sieve, self-interp meta-eval)
x 4 tiers (Python | C | Asm | All three)
- All-three mode renders the three tier outputs side by side with
per-tier elapsed timing
Tests (38 verified assertions)
- 20 unit (Node): per-tier module loads, eval smoke
- 8 integration (Node): each demo on c+asm WASM byte-matches the
canonical native Python run
- 10 functional (Playwright headless Chromium): page mounts, every
demo runs on every tier, all-three renders
Makefile
- Root targets: wasm-build, wasm-test, wasm-test-fn, wasm-serve,
wasm-deploy, wasm-clean
- wasm/Makefile orchestrates the three tier builds; deploy copies
dist/ into www/playground/
Asm tier notes
- WAT linear symbol intern + linear env lookup is MOAD-0001 at scale;
documented in the asm/lumbda.wat header and in the SPA footer. The
demos hit ~30 globals so the linear walks are cheap enough.
- Bump allocator never frees (matches asm/lumbda.s heap discipline);
memory.grow expands by 1 MB chunks. Browser tab tears down at unload.
Toolchain (developer prerequisites)
- Emscripten 6.0.0 via emsdk at ~/git/emsdk
- wabt 1.0.36 at ~/git/wabt
- Playwright for functional tests (symlinked from ~/git/agnt)
29 new files publish factory infra (V2 autoscaler with live VRAM
sampling + EWMA peak tracking, HUGE solo-dispatch, two-tier DLQ/rDLQ
classifier + retry), general quantum circuit primitives (Cuccaro
ripple-carry adder, Clifford gate library, Clifford tableau simulator,
mod-arith family, dialog GCD reversible inverse, Karatsuba multiplier,
Solinas fast reduction), and a TCRAUDT reducer harness. Originally
developed in ~/git/www.foxhop.net/ecdsa/ for secp256k1 attack-surface
research; published upstream as obligated by AGPLv3.
Parametrization contract at factory/CONTRACT.md. Consumers export
LUMBDA_REPO_DIR + LUMBDA_QUEUE_DIR + LUMBDA_BACKEND_CMD + LUMBDA_EMITTER_CMD
then exec factory scripts. No fork-and-modify; single source of truth
upstream.
Integration tests gate 7 V2 defect classes that wedged a live factory
on 2026-06-12 (skewed-demand starve, zero-floor reservation,
multi-tier greedy, +-25%% damping, cold-start ramp, DLQ surge halve,
post-damp CPU ceiling) + 28 DLQ classifier cases (auto-retry vs
escalate partition) + bash -n syntax lint across every script.
GPU backend stays in consumer trees; rationale in
factory/GPU-BACKEND-NOTE.md. Bend wire protocol + gpu-worker.lsp
already upstream at examples/cuda-fanout/.
make factory-lint bash -n on every factory/*.sh
make test-integration V2 reducer + DLQ classifier + syntax gate
make sweep-doctrine TCRAUDT reducer gate (serial)
make sweep-doctrine-parallel xargs -P fan-out
Verified on neoblanka: factory-lint 12 scripts PASS; test-integration
14 V2 cases + 28 DLQ classifier cases + 12 syntax cases all PASS.
Drop the static *vram-budget-mib*=22000 cap that fox flagged as wrong:
'we shouldn't limit with a max — the algo should determine how many
children based on the bend forms usage in vram.'
New algorithm:
- *gpu-total-mib* (24576 default, RTX 3090) + *gpu-headroom-mib* (1024 pad)
- *vram-per-cell-max-mib* (4096 seed) tracks largest cell observed.
- admit-fork? returns true iff
(current_vram + projected_cell + headroom) < gpu_total
- wait-admit blocks at run-loop top using projected = current per-cell
max. Self-tunes: tiny cells → many concurrent, huge cells → few.
Helper file-size-mib (stat -c %s) reads bin file size as cheap proxy
for per-cell VRAM (bin file on disk ≈ peak VRAM bend-cuda loads).
Open: cross-fork learning. record-cell-vram! runs IN THE CHILD so
parent's *vram-per-cell-max-mib* doesn't see updates without a fork-
shared signal (TODO: parent peek bin path before forking, or child
writes per-cell-size to small file the parent reads). For now the
seed value + max-tracking-in-future-runs handle the common case
where all cells are similar size.
Cross-tier API parity with c-tier (81ac49e) + python-tier — all three
lumbda runtimes now share the substrate for fork-per-accept patterns.
Implementation: direct syscalls (no libc):
- SYS_FORK=57 → bi_forkself, returns 0/pid via make_int
- SYS_WAIT4=61 + WNOHANG=1 → bi_waitpid_nonblock, returns pid or 0
- SYS_EXIT=60 → bi_exit_immediate (same as bi_exit on asm — no atexit
to bypass; present for cross-tier API parity)
- SYS_NANOSLEEP=35 → bi_sleep, stack-allocated timespec (tv_sec=N,
tv_nsec=0), returns VAL_VOID
Built + tested in vm-runner.sh VM (Ubuntu 2G/2vCPU): both lumbda
and lumbda-gc + fork-cycle test = 3/3 children reap clean, exit-
immediate returns to parent waitpid correctly. Same behavioral
contract as c-tier (commit 81ac49e) and python-tier.
Single-PID parent persistent listener; each accept forks a short-lived
child handler that owns one bend-cuda subprocess + responds + exits.
Linux COW handles memory; OS scheduler distributes across cores.
N concurrent requests = N children + parent — naturally VRAM-isolated.
VRAM admission: wait-vram-clear queries nvidia-smi before each fork,
blocks accept when used > *vram-budget-mib* (default 22000, override
via LUMBDA_VRAM_BUDGET_MIB env). 24G card with avg 1-2GB per bin
supports 4-12 concurrent comfortably.
Requires lumbda c-tier fork-self / waitpid-nonblock / exit-immediate /
sleep primitives (commit 81ac49e). Child uses exit-immediate not exit
to avoid dual-cleanup hang on shared parent state.
Substrate for fork-per-accept pattern in gpu-worker.lsp — enables
async bend dispatch with internal load balancing.
c-tier (builtins.c):
- bi_fork_self: fork() wrapper, returns 0 in child / pid in parent / #f on fail
- bi_waitpid_nonblock: waitpid(-1, WNOHANG), returns reaped pid or 0
- bi_exit_immediate: _exit() wrapper — REQUIRED in fork-self children,
regular exit() runs atexit handlers against shared parent state and
hangs the child (observed empirically 2026-06-11 via vm-runner.sh).
- bi_sleep: real wall-clock sleep(3) — yields CPU. Replaces busy-loop
patterns that would (a) burn CPU and (b) SIGKILL in cgroup-limited
VMs (observed: 100M iter let-loop SIGKILL'd after 5s in qemu vm).
python-tier (lumbda.py): _fork_self / _waitpid_nonblock / _exit_immediate
/ _sleep mirrors via os.fork / os.waitpid / os._exit / time.sleep.
Tested in vm-runner.sh VM (Ubuntu 2G/2vCPU): 3-child fork-cycle test
spawns + reaps cleanly 3/3 in both c-tier + python-tier. The exact
test pattern that crashed neoblanka host pre-fix now works fine.
Asm tier: deferred. Lock retained at chmod a-x ~/git/lumbda/asm/lumbda*
per CLAUDE.md threat model.
Earlier rename kept DEMO_OPS env fallback + *binary-demo-ops* var name
as transitional back-compat. With both hosts redeployed on bend-cuda
that's no longer needed.
Renamed:
*binary-demo-ops* -> *binary-bend-cuda*
DEMO_OPS env -> removed (only BEND_CUDA recognized now)
Also bulk-updated cuda-fanout sibling docs (DESIGN, CATALOG, plans/)
that still spelled the old name.
Slot reserved for future bend-rocm / bend-cpu via parallel env vars.
bend's (health) now reports whether $BEND_QUEUE_DIR/FEEDER_PAUSE
marker is set. Consumers (feeder, factory-status, ops scripts) get
pause state in the same single RPC as bend liveness + supervisor
proc counts + queue depths.
Returns:
1 marker present (operator wants this host out of rotation)
0 no marker (host in active rotation)
-1 BEND_QUEUE_DIR env unset (host has no associated queue)
Lets a future feeder version drop separate SSH pause-marker probes
in favor of the bend health RPC. Today's feeder still does the SSH
check; this just opens the door for the simpler model.
Mirror of pool-procs added earlier — bend health now reports both
pool-procs and dispatcher-procs counts so a single RPC tells the
feeder/factory consumer whether either supervisor is dead while
queue has work.
Failure mode this fixes: 4090's bend-dispatcher hit MAX_IDLE_LOOPS
drain-exit; balance moved 5 cells into its queue but no dispatcher
to consume them. Bend health response previously didn't surface the
gap; feeder's separate SSH probe (now added in www.foxhop.net commit)
caught it but a single RPC is cheaper than per-host SSH.
Follows foxhop ecdsa repo rename of infrastructure scripts. The pool
process is now named bend-emit-pool (job-agnostic), not ecdsa-emit-pool.
health-pool-procs pgrep updated to match.
bend co-lives with ecdsa-emit-pool on each foxhop production host.
When pool dies but bend stays up, .lsp cells pile un-emitted; bend
sits idle waiting for .ready bins that never arrive. Today's incident
took 30+ min to surface because feeder couldn't tell from bend health
alone — needed a separate SSH+pgrep per host.
Add pool/queue counts to (health) so one RPC returns the full picture:
(ok (load-avg L) (vram-free-mb V) (uptime-ms U)
(pool-procs P) (queue-ready R) (queue-emitting E) (queue-done D))
Helpers:
health-pool-procs pgrep -cf ecdsa-emit-pool
health-queue-count EXT ls $BEND_QUEUE_DIR/*.EXT | wc -l
BEND_QUEUE_DIR env var — set when bend is launched on a host with an
associated pool. Absent → queue counts return -1 (caller treats as
'unknown / not applicable').
Caller now has single-RPC view of bend + pool + queue health; feeder
can drop its separate SSH pool-watchdog probe in favor of the bend
(health) RPC field.
Lets bend (examples/cuda-fanout/gpu-worker.lsp) run multiple worker
processes behind a single listening port. Each worker calls
tcp-listen on the same port; the kernel distributes incoming
connections across the bound sockets.
Foxhop production use case: 2 bend workers per GPU host (3090 + 4090)
to consume the .ready queue at 2x throughput without an external
load balancer.
The CUDA binary that bend's gpu-worker.lsp spawns to run QECCOPS1
ops.bin sims is renamed from demo_ops to bend-cuda (production naming;
AMD ROCm equivalent will land as bend-rocm).
*binary-demo-ops* now reads BEND_CUDA env first, falls back to legacy
DEMO_OPS env, then default path /home/fox/git/www.foxhop.net/ecdsa/cuda/
bend-cuda. The legacy demo_ops symlink is in place at the install side
so back-compat holds during the rename transition.
See www.foxhop.net commit (sibling repo) for the foxhop ecdsa side
of the rename.
Root cause for repeated bend gpu-worker crashes mid-pipeline:
;;; bend RECV cuda-sim-ops-bin ops=monitor-r00007-c006.bin
error: not a pair: #f
When demo_ops crashes mid-write (GPU OOM, segfault, etc) the portal
file may exist but be empty or malformed. read-from-string returns
#f. Both portal-timing + portal-find-number then call (cdr #f) =
runtime crash, taking the port-8320 listener down.
Fix: gate both walkers on (pair? form) BEFORE descending. When form
is not a pair (= #f from empty/malformed portal), return #f cleanly.
Caller's (or cpu-ms 'NA) display already handles #f.
Watchdog still exists for true OOM / kill -9, but routine empty-portal
events no longer take bend down.
Reproduces by sending a malformed S-expr to portal-find-number or
portal-timing.
ecdsa's Phase B emit at full secp256k1 width writes a 4–6 GB QECCOPS1
ops binary. The existing path — accumulate in a string-output port,
materialize via get-output-string, write — peaks RAM at 3× body
size (port internal buffer + Scheme string copy + write-binary-file
concat). A 6 GB body needs ~18 GB transient; OOMs a 16 GB QEMU guest.
This commit shifts emit-stream onto a constant-RAM file-port path
and fixes binary-correctness defects in the supporting primitives.
New primitives (mirrored across c/builtins.c + lumbda.py):
- open-binary-output-file path
Opens in "w+b" so the caller can seek back to rewrite a header.
- port-set-position! port offset
fseek absolute offset on a file port. emit-stream reserves a
16-byte placeholder header, streams the body, then seeks back to
byte 0 to rewrite the QECCOPS1 + n_ops u64 LE once n_ops is known.
- append-binary-file path data
Opens in "ab" and fwrite's the bytes through. Pairs with
write-binary-file so callers can land header + body in two writes
instead of (string-append header body).
- append-port-to-binary-file path port
Streams a string-output port's buffer to disk via fwrite without
materializing (get-output-string port). Lets callers keep their
existing string-output sink and avoid the body-size string copy
if they stay on string-port emit.
Binary-correctness fixes:
- bi_write_string to a file port used fputs, which calls strlen.
Binary payloads containing 0x00 truncated at the first null byte.
Switched the file-port branch to fwrite with the string's known
->len (same fix family as the earlier bi_get_output_string
strlen defect).
- bi_write_char per-byte fflush guarded to stdout only. With
millions of gate-bytes per second, flushing after every fputc to
a file port was a 100× slowdown. File ports buffer until close
or explicit flush-port — keep stdout's per-byte feedback path,
drop fflush on every file-port byte.
- port_write_str grows 1.5× past 256 MB instead of 2× throughout.
At realloc time the transient peak is old + new; 2× at 8 GB →
16 GB transient needs 24 GB. 1.5× bounds peak at 2.5× and keeps
multi-GB string-port workloads inside a 16 GB VM.
Tests: 88/88 c-test, 4/4 regression-named-let-leak, 205/205
functional, zoe-favorites all tiers. Binary roundtrip with embedded
nulls at 10/1000/100000 bytes passes byte-for-byte.
End-to-end: foxhop ecdsa DIALOG_GCD secp256k1 emit lands a 4.7 GB
binary at 322 MB peak RSS in 7:41 wall on a 16 GB QEMU guest.
Two C-tier defects surfaced when running foxhop's ecdsa Phase B emit
through the lumbda C interpreter. Both produced wrong bytes in the
generated QECCOPS1 binary; both Python tier handled correctly.
1. string-ref on a byte >= 0x80 returned a char with codepoint -1.
The store was a `char` array (signed on x86_64); `s->data[idx]`
sign-extends 0xFF into a negative int before VAL_CHAR wraps it.
Round-tripping (char->integer (string-ref s 0)) for a 0xFF byte
gave -1 instead of 255. The Scheme-level (u64-le n) packer uses
(make-string 1 (integer->char (modulo v 256))) for each byte and
reads them back; sign-extension corrupted the high-bit bytes.
Cast through `unsigned char` in bi_string_ref.
2. pack_u64_slot treated bignum slots as raw fixnums.
ecdsa's emit-stream binds (no-slot *no-slot*) where *no-slot* is
18446744073709551615 (u64 max). On Python tier that's a regular
big-int. C tier carries it as a bignum NaN-box slot. The packer
knew about VAL_FALSE → 0xFFFF... but called as_int(v) on bignums,
reading the NaN-box payload bits (pointer-to-Bignum) and writing
that pointer as the field's 64-bit LE value.
Add an IS_BIGNUM branch that extracts the low 64 magnitude bits
directly. Cross-tier emit-stream code stays unchanged.
Tests: 88/88 c-test, 4/4 regression-named-let-leak, 205/205
functional. DIALOG_GCD smoke at p=11 n+1=5 now byte-identical to
Python tier (1,941,816 bytes match exactly).
Two pre-existing defects in c/builtins.c made (open-output-string)
unusable for binary emit:
- bi_get_output_string ran the buffer through make_string_from_cstr,
which calls strlen. Any 0x00 in the payload truncated body at
that byte. Use port's known str_len directly via make_string.
- bi_write_char ignored string-port destinations entirely — it
pulled AS_PORT(p)->fp (NULL for string ports), fell back to
stdout, and silently routed gate bytes to terminal output
instead of the port buffer. Route to port_write_str when target
port kind is PORT_STRING, matching bi_write_string's behavior.
Surfaced while validating the precise-GC fix against ecdsa's
Phase B emit (writes 56-byte op records full of embedded nulls
through a string-output port, get-output-string at the end).
With these fixes ecdsa's n+1=64 emit drops a 333 MB binary in
42.6 seconds — pre-fix it timed out at 600 s with a 17-byte
header-only file (strlen truncated body at byte 1; the visible
gate bytes had been escaping to stdout the whole time).
Binary roundtrip test (write n bytes alternating x / 0x00 to
output-string port, read back via get-output-string):
n=10 len=10 ok
n=1000 len=1000 ok
n=100000 len=100000 ok
All upstream tests still pass (88/88 c-test, 4/4 regression,
205/205 functional, zoe across tiers).
Boehm's conservative pointer scan cannot recognize lumbda's Value
layout — heap pointers live in the low 48 bits with QNAN + tag bits
in the upper mantissa, so a raw word never looks like a heap address.
Until now main.c neutralized this with GC_disable(): every allocation
leaked, OOMing any long-running workload.
Add precise tracing via a custom Boehm kind:
- New c/gc.c: mark proc walks 8-byte words in mixed mode — when the
QNAN bits are set with a pointer-bearing tag (0/2/4/5/6) extract
the low-48 pointer; otherwise fall through to raw-pointer
validation. GC_set_push_other_roots callback decodes NaN-boxed
Values on the C stack via setjmp anchor + scan up to the stack
base captured at process start.
- Allocations holding Values (Pair, Env bindings, ValueStack data,
ULVector data, HTEntry, Proc params + body, FullCont stack,
CodeObj instrs, SymbolEntry) route through lumbda_value_malloc.
Pure-byte sites (bignum limbs, char buffers, source files) stay
on regular GC_MALLOC.
- main.c / test.c / bench.c capture stack-base then drop GC_disable.
types.c also zeros popped slots on the value stack so stale pointers
do not survive a vs_pop and pin freed objects — independent
correctness fix that pays off once GC actually runs.
Build: USE_GC=1 (default when /usr/include/gc.h exists).
Tests with GC enabled:
- 88/88 c-test
- 4/4 regression-named-let-leak (test that motivated GC_disable)
- 205/205 functional (Python + C)
- zoe-favorites all tiers (Python + C + asm + asm-full)
alloc-test 1M cons drop-loop:
- Before: 0.60s wall, 156 MB RSS, leaks every cell
- After: 0.37s wall, 4 MB RSS, ~1500 GC cycles each freeing ~370 KB
Ports lumbda.py _emit_circuit_to_ops_bin_stream to C tier. Walks a
Scheme registers + ops list, writes each 56-byte QECCOPS1 op record
straight to disk via fopen/fwrite, then seeks back to patch the n_ops
header at the tail. O(1) host memory regardless of n_ops.
Cross-tier byte-identity verified:
* 500-op mixed-tag synthetic circuit on host: Python ↔ C identical
* real-point-add n+1=5 (p=11, 94,214 ops, 5.0 MB) inside VM
* real-point-add n+1=9 (p=251, 674,872 ops, 36 MB) inside VM
sha256 matches every width.
Speedup on equal Scheme source (build + emit combined):
* n+1=5: Python 56.8 s → C 1.1 s (52×)
* n+1=9: Python 365 s → C 8.4 s (43×)
Combined ratio exceeds the prior 11-16× C-tier envelope because the
build phase (Phase B mod-arith construction) also accelerates on C;
emit-only ratio is ~3-9× and grows with op count.
C test suite: 85 → 88 passing (emit-stream-basic, alloc-free, empty).
Shared functional suite: 205/205 still passing on both tiers.
Implementation notes:
* pack_op_record packs u32 kind + u32 pad + 6× u64 LE, matching
op-specs->bytes byte layout exactly.
* Layout hashtable Symbol → fixnum base, mirroring walk-circuit-ops.
* NO_SLOT sentinel: 0xFFFFFFFFFFFFFFFF written directly to u64 slots
that the Scheme side did not populate.
* libc's default fwrite buffer (~4 KB) handles batching at ~73 ops
per write — same throughput class as Python tier's 8 KiB list batch.
Folds 76 commits of substance since 2026-04-24 into the whitepaper
without losing any of the novel cross-domain glue the language earns
its keep on.
Additions
- §4.3 extended from four scopes to five, with bend as feedback across
heterogeneous compute (host ↔ GPU). New diagram
diagrams/five-scopes-feedback.png stacks the boundaries.
- §11.8 Bend: cross-tier GPU dispatch. Wire modes (S-exp text, BSHK
binary), worker hosting per tier, spawn-process-stdio + flush-port
cross-tier IPC, worker health heartbeat with VRAM-ranked pick,
6-row catalog of live forms (cuda-shake-fanout, cuda-secp256k1-mul,
cuda-bignum-cgbn, cuda-radix-sort, cuda-blake3-tree,
cuda-sim-ops-bin). New diagram diagrams/bend-dispatch.png.
- §2.1 C-tier bignum: arbitrary-precision integers (Boehm-GC managed),
which unblocked the secp256k1 widths the GPU forms need on the host
side.
- §12.1 post-cycle audit: recv-exact O(n²)→O(n), asm scheme_read
overflow, asm gc_sweep page-fault, C-tier JIT cur_code restore
across CALL/RETURN.
- §13 GPU Phase 1 moved from future to shipped; Phase 2 trampolining
and Phase 3 interaction combinators sharpened.
- Source repo link (git.unturf.com/engineering/unturf/lumbda) added
to the cover page band and the Citation block.
Compressions (no novel glue dropped — audited per fox's constraint)
- §8.2-§8.5 EML derivation prose collapsed into a single §8.2
derivation chain code block; §8.6 renumbered to §8.3.
- §6.6.3 adaptive-meta-GC narrative compressed; result table kept.
- §7.4.1 GC-build S-expression portal tightened to two paragraphs;
"language is its own wire format" insight kept.
Preserved in full per audit
- §7.5 portable RNG state across tiers
- §7.5.1 one-side kernel entropy + portable bit-identical continuation
- §11.7 sendfile + adaptive preload (27 KB asm-gc within 2% of Caddy)
- §6.6.1 collaborative arena + mark verifier
- §6.6.5 precise block typing
- §8.1 + §8.3 Lumbda hosts its own EML proof checker, ~16× faster
than Lean cold
Net RST: 1622 → 1655 lines.
inject-whitepaper-css.py adds CSS for a fixed 280px left sidebar plus a
DOMContentLoaded script that walks `section[id] > h2, > h3` and builds a
chapter/section nav. IntersectionObserver tracks active section as you
scroll; the rail scrolls itself to keep the current entry in view. Below
1100px the rail hides behind a hamburger toggle with backdrop dismiss.
Cancels docutils responsive.css `body > *` blanket padding on the rail
via the same `all: revert` pattern already used for the uncloseai
floating button. 59/59 source sections present in rendered DOM.
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.
Adds tagged bignum support alongside the existing 48-bit fixnum on the C
tier. Tag 6 = bignum, heap struct sign-magnitude with u64 little-endian
limbs. Reader emits bignums for any literal past the fixnum range; +, -,
*, quotient, remainder, modulo, expt, =, <, >, abs, odd?, even?,
integer?, exact?, number->string, string->number all promote fixnum →
bignum on overflow & demote back when results fit. Boehm GC owns every
allocation. Schoolbook O(n²) mul + shift-subtract divmod is sufficient
at our 4-limb / 256-bit scale.
Before: (expt 2 48) = 0, (expt 2 256) = 0, secp256k1-p = -4294968273.
After: all three return their exact arbitrary-precision values, matching
Python tier byte-for-byte.
Validated:
- c/test.c — 85/85 pass (+2 new bignum unit tests).
- tests/functional.lsp — 205/205 pass on both C & Python tiers.
- tests/bignum-cross-tier.lsp — 33/33 pass byte-identical on both tiers
(diff produces no output).
- ecdsa/runs/lumbda-sweep-003/c-tier-bignum-probe.lsp — all four
assertions now match the Python oracle.
- ecdsa Phase B byte-identity sweep inside QEMU guest:
n+1=9 p=251 sha256 c668bbe3... — matches Python oracle.
n+1=18 p=131071 sha256 8a031f96... — matches Python oracle.
n+1=33 p=2³²-5 sha256 0bc56905... — matches Python oracle.
Previously the n+1=33 C tier emitted sha256 b024d6d9... (26,078 fewer
Toffolis due to silent fixnum wrap). Bignums close that gate.
secp256k1 production-width emit (n+1=257) is now structurally unblocked
on C tier; downstream agent (#55) drives that next-step on the ecdsa
side. Asm tier inherits in a follow-up port.
Foxhop ecdsa point-add emit at n+1=257 secp256k1 (~15M ops) overruns
host RAM via the accumulator path — walk-circuit-ops builds a 4 GB
Scheme list, op-specs->bytes builds 840 MB body string before
write-binary-file ships it. sweep-secp256k1.lsp OOMs at 8 GB.
Streaming primitive walks the circuit & writes each 56-byte op record
directly into an open binary file handle in a single pass:
(emit-circuit-to-ops-bin-stream out-path registers ops)
Per-op: pack 56 bytes via the same struct.pack format as
op-specs->bytes, append to an 8 KiB batch buffer, flush to disk when
threshold hits. After the walk closes, seek(8) patches the n_ops
header field. No intermediate list, no intermediate body bytes.
Memory profile at n+1=32 / 146 ops: ~64 KB grew during emit (host
file buffer + interpreter overhead). Holds flat at n+1=256 / 1154 ops
& n+1=1024 / 4610 ops. Will hold flat at n+1=257 secp256k1 / ~15M
ops — the byte-level operation count grows linearly but RSS does not.
Dispatch table mirrors _walk_circuit_ops exactly so bytes are
bit-identical to (op-specs->bytes (walk-circuit-ops ...)) by
construction. Verified byte-for-byte at p=11 (48 763 ops), p=251
(159 263 ops), & a hand-rolled n+1=32 hand circuit.
Foxhop wrapper lands as (emit-ops-bin-stream out-path c) in
ecdsa/lumbda/emit-ops-bin.lsp. Existing (emit-ops-bin ops out-path)
accumulator API stays untouched for backward compatibility — all
288 ecdsa/tests/unit/test-emit-ops-bin.lsp assertions still pass.
Earlier port-flip commit (8d66bc0) accidentally added a submodule
pointer at .claude/worktrees/agent-a7b0eac1 via a stale entry in
lumbda.py's git add -A run. No .gitmodules file ever existed for it,
so clones would carry a dangling gitlink. Remove from index & add
.claude/ to .gitignore so agent state stays untracked.
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).
Direct C translations of our Python-tier primitives at lumbda.py
_walk_circuit_ops / _op_specs_to_bytes / _count_lumbda_ops (commit
99701c8). Same algorithms; native dispatch via interned-Value identity
on cached symbol globals.
NO_SLOT representation differs from Python tier. Python stores the
literal 18446744073709551615 fixnum (arbitrary-precision int). C-tier
fixnums cap at 48 bits via NaN-boxing PAYLOAD_MASK, so we substitute
VAL_FALSE as our slot sentinel inside op-spec vectors; pack_u64_slot
writes 0xFFFFFFFFFFFFFFFF whenever it sees VAL_FALSE. Both tiers
produce byte-identical QECCOPS1 output.
API:
- (walk-circuit-ops registers ops) -> list of 7-element op-spec vectors
- (op-specs->bytes specs) -> latin-1 string of 56*N bytes
- (count-lumbda-ops ops) -> 3-element vector (toffoli clifford total)
Symbol cache (g_sym_ccx / g_sym_x / ...) initializes lazily on first
call; intern() is idempotent so repeat init costs nothing. Layout uses
make_hashtable + ht_set/ht_ref/ht_delete for O(1) qubit-base lookup
matching our Python dict-based implementation.
Measured wall on foxhop ecdsa's canonical p=11 textbook+refined emit
pair inside our QEMU guest:
Python tier: 6.9 sec (baseline after primitive lift)
C tier: 0.71 sec (9.7x over Python)
Output verified byte-identical against Python-tier reference files via
cmp on both textbook & refined paths.
Asm tier port deferred. Asm tier's documented role serves bulk 9024-
shot validation (simulator runs against an emitted ops.bin, 160x Python
on portal round-trip per whitepaper s6.6.4) — emit pipeline targets
Python / C tier. ~700-900 lines of hand-written x86_64 asm + QEMU
debug cycles, 2-4 day effort, no current asm-tier emit consumer.
Adds a (health) op handler on the worker side & a lazy-refresh
health cache + VRAM-aware selection on the bend client side.
WORKER (gpu-worker.lsp)
(health) returns (ok (load-avg L) (vram-free-mb V) (uptime-ms U))
- L from /proc/loadavg first field
- V from `nvidia-smi --query-gpu=memory.free --format=csv,noheader,nounits`
(returns 0 when nvidia-smi missing — host w/o NVIDIA GPU)
- U from current-time-ms; client detects a worker that hung
& restarted between probes via uptime jump
Backward-compat: workers without (health) return
(error (unknown-op health)); client treats that as ok+vram=0.
CLIENT (bend.lsp)
*worker-health* alist keyed "host:port" → (last-checked-ms status vram-mb)
Cache TTL on ok = 5 s; cooldown on down = 30 s.
bend-pick-worker now:
- filters out workers in down-cooldown
- sorts healthy peers by free VRAM descending
- falls back to round-robin if every worker is in cooldown
bend-dispatch-to-gpu flips workers to down on tcp-connect-fail
or empty-reply so a transient failure costs at most one call.
Two lumbda quirks caught while building:
- (eq? 0 #f) → #t in lumbda. worker-probe-health returns 0
(a number) for the unknown-op fallback, but if we'd checked
(eq? vram #f) we would have mis-marked the worker down.
Now uses (number? vram) instead.
- tcp-connect raises a Python ConnectionRefusedError (NOT a
LispErr) on dead-host probes. lumbda's `guard` only catches
LispErr; only with-exception-handler catches Python
exceptions. Probe now wraps via with-exception-handler so a
single dead worker never aborts a fleet iteration.
Smoke on Python tier:
mixed (127.0.0.1:1 dead + 3090-ai live) → cache shows down for
the dead one (30s cooldown), ok for live (22777 MB free VRAM,
measured by the worker's nvidia-smi probe).
Three Python-tier primitives that lift foxhop ecdsa's emit-ops-bin
pipeline out of the Scheme interpreter:
- walk-circuit-ops registers ops → list of op-spec vectors. Mirrors
ecdsa/lumbda/emit-ops-bin.lsp's walk-op dispatch table (alloc, free,
x, z, cx, cz, swap, ccx, ccz) in pure Python with interned-symbol
identity dispatch & a Python dict for the qubit-layout. Replaces a
Scheme named-let walk that paid ~5-9 ms per op via per-iteration
closure / let* / cons / append overhead — wall dropped 252 sec to
237 ms on a 32 k-op p=11 case.
- op-specs->bytes vector-list → 56*N latin-1 string. Each op-spec is a
7-element vector packed via struct.Struct('<IIQQQQQQ').pack; results
joined once with b''.join + .decode('latin-1') so the existing
write-binary-file primitive ships the body byte-for-byte. Replaces
per-op (string-append (u32-le ...) (u32-le 0) (u64-le ...) ...) in
Scheme that paid ~2-3 ms per op; wall dropped 49 sec to 84 ms on the
same case.
- count-lumbda-ops ops → vector(toffoli clifford total). Tags ccx →
toffoli, x|cx → clifford, anything else → total only. Replaces a
pure-Scheme named-let count that hit ~200 sec per variant on a 32 k-op
list.
Net foxhop ecdsa wall on the canonical p=11 textbook+refined emit pair
dropped from ~10 min to ~7 sec — ~85x end-to-end on Python tier. Output
byte-identical against a pre-rewrite reference file (cmp clean on both
textbook & refined paths).
C-tier port: TBD (same algorithms, separate translation unit).
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.
Built op_id 0x0A (mod-inv-batch) on cgbn-batch-worker.cu using
Montgomery's 1 inv + 3(N-1) muls trick. Byte-identity holds at
every N (n in {32, 1k, 10k, 100k}, all 10 ops × 4 N = 40 PASS).
Kernel-ms 0x05 (per-instance, current LIVE op) vs 0x0A on a 3090:
N 0x05 0x0A ratio
10k 0.52 ms 78.34 ms 0.01x
100k 2.54 ms 762.28 ms 0.003x
1M 21.10 ms 7763.45 ms 0.0003x
The spec premise — "sequential modmuls cheap, parallel mod-invs
expensive" — INVERTS on a 3090. CGBN's parallel 0x05 saturates 82
SMs × 128 in-flight instances at ~47 Mops/s; one TPI=8 lockstep
instance walking 3N sequential modmuls in Phase 1+3 is
latency-bound, not throughput-bound. The Montgomery trick only
wins on hardware where one inv is dramatically more expensive
than 3N muls; on a 3090 the parallelism budget makes the
inversions cheap.
HARD-RULE TRIPPED (>= 5x win at n=100k required). Hard rule honored:
- no master daemon restart
- no CATALOG.md / bend.html promotion
- production daemons on :9091 / :9092 untouched (still serve 9-op binary)
Code + test changes & progress doc commit here as research artifact.
The 10-op binary builds clean & is byte-correct; just slow. Future
day-2 pass should land per-block parallel prefix scan (Kogge-Stone
or Sklansky) for Phase 1+3 — same refactor pattern Form A Day-4
needs to make v4 beat v3.
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.
handle-binary-shake, handle-binary-cgbn, handle-binary-secp each
called (delete-file in-path) (delete-file out-path) unconditionally
in both the daemon-ok and daemon-error branches — 6 sites total.
When the daemon failed without producing out-path (e.g. crash, OOM,
bad payload), delete-file raised file-not-found and the entire
listener exited. Crashed 3090-ai once during ai.foxhop.net deployment
smoke.
Wrap every delete-file with (if (file-exists? PATH) (delete-file PATH))
so a missing portal cannot kill the worker. (ok pong) sanity check
passes on Python tier.
Sites (12 guards = 6 path pairs × 2 paths):
handle-binary-shake : ok branch + error branch
handle-binary-cgbn : ok branch + error branch
handle-binary-secp : ok branch + error branch
