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
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.
Adds *bend-workers* list with round-robin dispatch, BEND_WORKERS env
loader, and helpers (bend-set-workers!, bend-pick-worker,
bend-parse-workers-env, bend-load-workers-from-env!). Single-host
legacy callers unaffected — when *bend-workers* is empty the
dispatcher falls back to *bend-worker-host* / *bend-worker-port*.
Validated cross-tier (Python + C lumbda) against a live two-host
cluster (3090-ai.foxhop.net:9091, ai.foxhop.net:9092) — round-robin
distributes evenly; cgbn mod-mul results byte-identical to gmpy2
reference on both hosts.
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).
Per Day-2 progress doc the kernel bottleneck was scalar_mul (256 doubles +
~128 adds per scalar), not _ModInv. v3 swaps the binary double-and-add for
a windowed-base ladder: precompute table[0..15] = i*P affine on-device,
walk scalar 4 bits at a time MSB->LSB, cutting per-scalar adds from ~128
to ~63. Table loaded into __shared__ (1024 B) per block.
Benched on 3090-ai (best-of-3, --no-batch-inv):
n=10k: v1 4.22 Mkeys/s v3 2.18 Mkeys/s 0.52x (init overhead dominates)
n=100k: v1 7.48 Mkeys/s v3 9.29 Mkeys/s 1.24x
n=1M: v1 7.87 Mkeys/s v3 13.60 Mkeys/s 1.73x
Byte-identical against coincurve at n in {32, 1000, 10000, 100000, 1000000}.
CLI: --window-w 4 selects v3 (Day-3 canonical). w=8 reserved but stub-
rejected since device-side table init is register-stack-bounded at W <= 16.
Default no-flag behaviour stays v1 (Day-1) so the gpu-worker.lsp daemon
inherits the safe baseline until fox routes traffic to v3.
Companion progress doc at plans/form-A-day3-progress.md.
Captures empirical record from three agents that finished today:
form-A-day2-progress.md (STOPPED, no merge)
v2 Montgomery batch inversion regressed v1 by 0.63x-0.94x
across n in {10k, 100k, 1M}. Byte-identity PASS at every N;
math correct. Root cause: at this N, scalar_mul (256 Jacobian
doubles x 5-6 ModMult each) dominates, NOT _ModInv. v2's Phase
B/D used 1 thread/block leaving ~97% of SMs idle. v1 baseline
re-measured at 7.88 Mkeys/s at n=1M (catalog upward correction
from initial 6.51). Day-3 path: warp-level prefix scan OR
windowed-G ladder.
form-D-axis-flip-RESULTS.md (Form D opt 1 shipped foxhop 1f7ac9d)
Per-candidate axis kernel measured 217 Mops/s at K=32 M=4 on
3090; 23.7x over per-shot N=4 at same M. Both axes saturate
at the same ~220-250 Mops/s, refuting the bandwidth-bound
diagnosis. Axis flip's win is occupancy-amortization at small
M, not bandwidth redistribution.
form-D-build-progress.md (Form D AG kernel parked)
Aaronson-Gottesman tableau dead-end: our point-add circuit
contains no H or S, so state never leaves the computational
basis and AG buys nothing. Three pivot options proposed; fox
picked options 1 + 3 in parallel.
Day-1 binary at lumbda 7661788 stays canonical for
cuda-secp256k1-batched-mul; v2 working-tree code lives uncommitted
on the build host as Day-3 scratch.
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).
QECCOPS2 packed-op format landed in foxhop ecdsa/cuda/ at commit
90484ca. 28 → 24 B per op in VRAM, 56 → 24 B on disk (2.33× shrink).
n_batches 16/64/128 on RTX 3090: ~7% kernel speedup, byte-identical
CPU vs unpacked-GPU vs packed-GPU. Bandwidth-bound diagnosis from
form-D-build-progress.md §4 stands — per-shot state traffic (~85×
larger than op stream) owns the 1.07× ceiling. Per-candidate axis-
flip (sim_gpu_axis.cu, foxhop 1f7ac9d) remains the open lever; this
~7% stacks on top.
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.
Per examples/cuda-fanout/plans/form-A-secp256k1-batched-mul.md.
Batched secp256k1 scalar*G via per-thread Jacobian double-and-add
plus per-thread Z-inversion to affine. Field arithmetic uses
FixedPaul/VanitySearch-Bitcrack GPUMath.h verbatim
(commit 66e6f9d, AGPL-3.0, vendored under
vendor/vanity-search-bitcrack/).
Wire: BSCP request / BSCR response distinct from BSHK/BCGB.
"BSCP" u32 op_id u32 n base_xy(64B BE) scalars(n*32B BE)
"BSCR" u32 status u32 n points(n*64B BE x||y)
Validation against coincurve on 3090-ai.foxhop.net byte-identical
across known-small (k in {1,2,3,7,0xdeadbeef,n-1,n,2^128-1})
& random sweeps at n in {32, 1k, 10k, 100k}.
Measured throughput on 3090, kernel-only:
n=10k 2.32 ms 4.31 Mkeys/s
n=100k 15.37 ms 6.51 Mkeys/s
End-to-end over warm TCP daemon from another host:
n=100k 400 ms 250 kkeys/s wall (PCIe + wire serialization
bound; kernel still <16 ms)
Speedup vs coincurve CPU single-thread (~44 kkeys/s host)
~148x at n=100k kernel-only. Day-2 work to add _ModInvGrouped
batched inversion should push toward FixedPaul's 6.9 Gkeys/s
published on 4090.
gpu-worker.lsp: maybe-register-daemon! for cuda-secp256k1-batched-mul,
handle-binary-secp branch in handle-one dispatching on BSCP magic.
Makefile: secp256k1-batch-mul / secp256k1-test / secp256k1-bench
targets. Test harness ships with coincurve preferred, falls back to
python-ecdsa or pure-Python double-and-add for the host oracle.
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.
