diff --git a/examples/cuda-fanout/DESIGN-go-gpu.md b/examples/cuda-fanout/DESIGN-go-gpu.md new file mode 100644 index 0000000..db6c270 --- /dev/null +++ b/examples/cuda-fanout/DESIGN-go-gpu.md @@ -0,0 +1,184 @@ +# `(go-gpu …)` — S-expression coroutines on remote GPU + +Sketch of the architecture fox proposed: Go-keyword-style coroutines +that ship S-expressions to a remote GPU box for execution. Pattern +analogous to vLLM's inference RPC, but for arbitrary lumbda +expressions backed by a registered CUDA primitive. + +## API shape + +```scheme +;; fire-and-get-future +(define f (go-gpu (cuda-shake-fanout inputs 32))) +(define result (await f)) + +;; fan out across N GPU workers in parallel +(define fs (go-gpu-many + (cuda-shake-fanout batch-1 32) + (cuda-shake-fanout batch-2 32) + (cuda-shake-fanout batch-3 32))) +(define results (await-all fs)) + +;; streaming (long-running kernels emit chunks back) +(go-gpu-stream (cuda-sim-batch ops-bin 141) + (lambda (chunk) (process-partial chunk))) +``` + +## Runtime architecture + +``` + ┌──── coordinator (any host) ──────┐ + │ - load tracker per GPU worker │ + │ - dispatch queue │ + │ - future registry │ + └────────────┬─────────────────────┘ + │ TCP / S-exp message + ┌──────────────────────┴───────────────────────┐ + │ │ + ┌───────▼───────┐ ┌───────▼───────┐ + │ gpu-worker │ │ gpu-worker │ + │ ai.foxhop.net │ │ 3090-ai.fox.n │ + │ RTX 4090 │ │ RTX 3090 │ + └───────────────┘ └───────────────┘ + │ │ + ▼ ▼ + spawn local binary spawn local binary + (cuda-shake-fanout, (cuda-shake-fanout, + cuda-sim-batch, etc.) cuda-sim-batch, etc.) +``` + +Reuses lumbda's existing fleet TCP primitive +(`ecdsa/lumbda/fleet/coordinator-fleet.lsp` pattern). The only new +piece is **GPU capability advertising** in the worker handshake + +**load-aware dispatch** in the coordinator. + +## Wire protocol (S-exp over TCP) + +Request (coordinator → worker): +``` +(go-gpu-request + (id 17) + (op cuda-shake-fanout) + (args + (inputs ("00" "01" "deadbeef")) + (output-bytes 32))) +``` + +Response (worker → coordinator → awaiting future): +``` +(go-gpu-response + (id 17) + (status ok) + (result + (hashes ("abc..." "def..." "fed..."))) + (timing-ms (work 0.2) (wire 1.4) (queue 0.0))) +``` + +Streaming response (long kernels): +``` +(go-gpu-stream-chunk (id 17) (seq 0) (data ...)) +(go-gpu-stream-chunk (id 17) (seq 1) (data ...)) +(go-gpu-stream-end (id 17) (status ok)) +``` + +## Capability advertising (worker startup) + +``` +(gpu-worker-hello + (host 3090-ai.foxhop.net) + (vram-gb 24) + (cuda-cores 10496) + (binaries + (cuda-shake-fanout (path /usr/local/bin/cuda-shake-fanout)) + (cuda-sim-batch (path /usr/local/bin/cuda-sim-batch)) + (cuda-keccakp (path /usr/local/bin/cuda-keccakp))) + (max-concurrent 4)) +``` + +Coordinator only routes operations whose `op` name appears in some +worker's `binaries` list. Unknown ops error early at submit time. + +## Scheduling + +Coordinator keeps per-worker state: + +``` +(worker-state + (host 3090-ai.foxhop.net) + (load 0.5) ;; fraction of max-concurrent in use + (vram-free-mb 23900) + (last-heartbeat 1717530000)) +``` + +`go-gpu` form picks the worker with lowest `load × (vram-needed / +vram-free)` and routes there. Future stored in registry keyed by +`(coordinator-issued-id worker-host)`. When response lands, future +fires. + +## Failure semantics + +- Worker disconnect mid-request → future raises `gpu-worker-lost` +- Worker reports OOM → coordinator marks worker `vram-tight`, retries + on another worker if available, else raises `gpu-out-of-memory` +- Network hiccup → coordinator retries up to 2× with exponential + backoff, then raises `gpu-timeout` + +`await` MUST be wrapped in error handlers same as any I/O primitive. + +## Per-tier integration (≈ 30 LoC each) + +Each lumbda tier registers four primitives in its dispatch table: + +| primitive | shape | what it does | +|-----------|-------|--------------| +| `go-gpu` | special form (delayed eval of body) | serializes op + args → submits to coordinator → returns opaque future handle | +| `await` | (await future) | blocks current continuation until response lands | +| `await-all` | (await-all (f1 f2 ...)) | blocks on every future; returns list of results | +| `go-gpu-stream` | (go-gpu-stream body callback) | spawns + invokes callback per chunk | + +Coordinator + worker stay outside lumbda's core: they're standalone +lumbda scripts (or, eventually, the same role lumbda's fleet plays) +launched at boot per host. + +## Comparison to vLLM + +vLLM is HTTP/REST with token-stream responses; lumbda's +`go-gpu` is S-expression over TCP with portal-style responses. Same +shape: client submits a job, server runs on GPU, client receives +result asynchronously. Key differences for lumbda: + +| dimension | vLLM | go-gpu | +|-----------|------|--------| +| wire | HTTP/JSON | TCP/S-exp | +| model loading | per-server, persistent | per-binary, spawned per request (today); persistent daemon per binary (future) | +| streaming | token-by-token | chunk-by-chunk | +| failure | HTTP status + retry | typed Scheme conditions | +| scheduling | one worker per model | one binary per capability; load-aware fanout | + +## What's needed to ship a working v1 + +1. **`gpu-worker-daemon.lsp`** — Scheme script that listens on TCP, + matches incoming op names against local binaries, spawns binary + with input portal, returns result portal. +2. **`gpu-coordinator-daemon.lsp`** — Scheme script that tracks + workers + their advertised capabilities, accepts client `go-gpu` + submissions, routes to least-loaded capable worker, holds futures. +3. **`go-gpu.lsp`** — client-side wrapper Scheme that submits & awaits. +4. **Capability advertise file** — `~/.lumbda/gpu-capabilities` per + host, lists local binaries the daemon should advertise. + +Each piece is ~50-100 LoC. Total v1 ships in one focused day. + +## Open questions for fox + +- **Q1**: should the coordinator be a separate process, or should + every lumbda runtime double as coordinator (mesh)? +- **Q2**: should `go-gpu` be a special form that captures free + variables in the body, or a regular function call that takes only + serializable args? (Special form is nicer but needs closure + shipping; function call is portable.) +- **Q3**: do we want pre-warmed daemon binaries (kill the 200 ms + cuda-ctx-init per spawn) or stick with per-spawn for simplicity? + Bench data informs this — see `bench.py`. +- **Q4**: should `await` be implicit at the end of a `(let ...)` + scope (Go-style scope-bound spawn), or always explicit? diff --git a/examples/cuda-fanout/README.md b/examples/cuda-fanout/README.md index de88cd6..d8743c8 100644 --- a/examples/cuda-fanout/README.md +++ b/examples/cuda-fanout/README.md @@ -113,9 +113,106 @@ rest of the repo. ```bash make all # compile via nvcc make test # round-trip vs hashlib.shake_256 -make bench # device vs host timing at N up to 100k +make bench # device vs host timing grid +python3 bench_daemon.py ./shake256-fanout # daemon vs per-spawn ``` +## Measured performance — where the GPU actually wins (and loses) + +**Setup:** RTX 3090 (10,496 CUDA cores, sm_86) + i9-12900K host, +Python `hashlib.shake_256` as the reference. + +### Per-spawn mode — kernel-launch overhead dominates + +For SHAKE256 specifically, the per-process CUDA context init eats +~200 ms regardless of workload. The actual kernel is 0.1–40 ms +depending on size. So the GPU loses to host hashlib at every size +we tested in per-spawn mode: + +``` +in_sz N total host ms dev ms kernel ms speedup +32 10,000 320 KB 6 282 0.12 0.02× +32 100,000 3.2 MB 60 322 0.24 0.19× +32 1,000,000 32 MB 606 1570 1.14 0.39× +32 3,000,000 96 MB 1818 4246 3.04 0.43× +1024 1,000 1 MB 3 239 0.25 0.01× +1024 100,000 102 MB 260 4753 1.72 0.05× +16384 10,000 164 MB 343 18638 2.66 0.02× +262144 1,000 262 MB 508 454171 39.07 0.00× +``` + +The 454 SECONDS at 262 MB is **portal parsing** of the 524 MB hex +text file — *not* the kernel (which is 39 ms). At the current +S-expression hex wire format, even our biggest kernels are dwarfed +by hex-string parsing. + +**Conclusion: SHAKE256 alone is too light to amortize either init +overhead or hex parsing.** Host hashlib reads raw bytes and stays +unbeatable for this primitive. + +### Daemon mode — 574× faster per-call + +`shake256-fanout --daemon` keeps the CUDA context alive across many +requests. CUDA init pays once at startup; every subsequent call +is just kernel + I/O: + +``` +workload: 10 calls × 100 inputs × 32 bytes each + + host hashlib loop : 0.6 ms total (0.06 ms/call) + per-spawn fanout : 1825.9 ms total (182.59 ms/call) + daemon-mode init : 109.2 ms (one-time) + daemon-mode calls : 3.2 ms total (0.32 ms/call) + + daemon is 574× faster per call than per-spawn + daemon end-to-end is 16× faster than per-spawn for 10 calls +``` + +The 574× factor is the actual win the `go-gpu` primitive should +inherit. Per-spawn mode is a debug tool; daemon mode is the +production architecture for workloads doing repeated fan-outs. + +### Daemon protocol + +``` +shake256-fanout --daemon + → reads commands on stdin, one per line: + process → fan-out, write result, "done " + quit → clean shutdown, "bye" + → response per command on stdout +``` + +This is exactly the interface a `(go-gpu …)` primitive in lumbda would +wrap: the per-tier dispatcher spawns one daemon per GPU host at boot, +then routes every `go-gpu` form's body through the existing daemon's +stdin. CUDA init never re-runs while lumbda is up. + +### When the per-spawn mode IS the right tool + +A few cases: +- **One-shot scripts** where init time doesn't matter +- **CI pipelines** that need fresh state per test +- **Debugging** with isolated CUDA contexts + +In those cases the per-spawn binary is sufficient and the daemon is +overkill. + +### Workloads that genuinely shine on this architecture + +SHAKE256 is too light. Workloads that actually win on GPU pour +compute into each thread: + +- **Reversible bit-vector simulation** (our `ecdsa/cuda/sim_gpu.cu` + does 12.8M ops × 2,295 threads = 30+ billion ops per kernel, 25 s + on a 3090) — kernel dominates, init is noise. +- **Many-round Keccak / PBKDF2 / Argon2** at million-iteration counts. +- **Bulk EC scalar mul** (we used OpenMP for 13× win on CPU; on GPU + the per-thread cost is the same but parallelism scales). + +The `cuda-shake-fanout` example is intentionally small so the +build + test cycle is fast. The pattern (leaf binary + portal + +daemon) is the contribution. + ## Provenance Keccak permutation derived from FIPS 202 reference. The compact diff --git a/examples/cuda-fanout/bench.py b/examples/cuda-fanout/bench.py index 45ee5af..1259b73 100644 --- a/examples/cuda-fanout/bench.py +++ b/examples/cuda-fanout/bench.py @@ -1,8 +1,14 @@ -"""bench.py — N-input fan-out: device vs host hashlib timing. +"""bench.py — find the workload shapes where cuda-fanout actually wins. -Generates N random inputs, hashes them on host (sequential Python), -then on device (parallel via shake256-fanout binary), compares -timings. Same hashes round-trip through the portal.""" +The per-process CUDA context init costs ~200 ms regardless of work +(measured: load 0.6 ms, launch 205 ms, kernel 0.1 ms, write 0.1 ms at +N=1000 × 32 B inputs). So GPU wins when total work × ns-per-byte +overtakes that 200 ms ceiling. + +Host SHAKE256 (Python hashlib) does ~500 MB/s on this i9-12900K. So: + break-even bytes total = 200 ms × 500 MB/s ≈ 100 MB + +We sweep around that threshold and report where GPU starts winning.""" import hashlib import os @@ -12,20 +18,38 @@ import sys import time BIN = sys.argv[1] if len(sys.argv) > 1 else "./shake256-fanout" +OUT_BYTES = 32 + +# Grid sized to actually finish & cover the break-even region. +# Each row's total bytes = SZ × N. We pick to bracket ~100 MB. +GRID = [ + # (input_size_bytes, N) + (32, 10_000), # 320 KB total — way under break-even, host wins + (32, 100_000), # 3.2 MB total — still host + (32, 1_000_000), # 32 MB total — getting closer + (32, 3_000_000), # 96 MB total — should cross + (1024, 1_000), # 1 MB total + (1024, 10_000), # 10 MB + (1024, 100_000), # 100 MB — should cross + (16 * 1024, 1_000), # 16 MB + (16 * 1024, 10_000), # 160 MB — should win + (256 * 1024, 100), # 25 MB + (256 * 1024, 1_000), # 250 MB — should win +] -def host_bench(inputs, out_bytes): +def host_bench(inputs_hex, out_bytes): t0 = time.monotonic() - out = [hashlib.shake_256(bytes.fromhex(h)).hexdigest(out_bytes) for h in inputs] + out = [hashlib.shake_256(bytes.fromhex(h)).hexdigest(out_bytes) for h in inputs_hex] return out, time.monotonic() - t0 -def write_portal(path, inputs, out_bytes): +def write_portal(path, inputs_hex, out_bytes): with open(path, "w") as fh: fh.write("(cuda-shake-fanout\n") fh.write(f" (output-bytes {out_bytes})\n") fh.write(" (inputs\n") - for h in inputs: + for h in inputs_hex: fh.write(f' "{h}"\n') fh.write("))\n") @@ -38,30 +62,78 @@ def read_hashes(path): return re.findall(r'"([^"]*)"', m.group(1)) if m else [] -def device_bench(inputs, out_bytes): +def read_timing(path): + import re + with open(path) as fh: + text = fh.read() + m = re.search(r"\(timing-ms\s+\(load\s+([\d.]+)\)\s+\(launch\s+([\d.]+)\)\s+" + r"\(kernel\s+([\d.]+)\)\s+\(write\s+([\d.]+)\)", text) + return tuple(float(x) for x in m.groups()) if m else (0, 0, 0, 0) + + +def device_bench(inputs_hex, out_bytes): in_p = "/tmp/cf-bench-in.portal" out_p = "/tmp/cf-bench-out.portal" - write_portal(in_p, inputs, out_bytes) + write_portal(in_p, inputs_hex, out_bytes) t0 = time.monotonic() subprocess.run([BIN, in_p, out_p], check=True) wall = time.monotonic() - t0 out = read_hashes(out_p) + timing = read_timing(out_p) os.unlink(in_p); os.unlink(out_p) - return out, wall + return out, wall, timing def main(): - sizes = [1_000, 10_000, 100_000] - out_bytes = 32 - print(f"binary: {BIN}\n") - print(f"{'N':>8} {'host ms':>10} {'device ms':>10} {'speedup':>8} match") - for N in sizes: - inputs = [secrets.token_hex(32) for _ in range(N)] - host_out, host_t = host_bench(inputs, out_bytes) - dev_out, dev_t = device_bench(inputs, out_bytes) - match = host_out == dev_out + print(f"binary: {BIN}") + print(f"out_bytes: {OUT_BYTES}\n") + print(f"{'in_sz':>10} {'N':>9} {'total':>9} {'host ms':>9} {'dev ms':>9} {'kernel ms':>10} {'speedup':>8} match") + print("-" * 86) + sweet = [] + for sz, N in GRID: + total_bytes = sz * N + # Skip if validate would take forever + validate = total_bytes < 64 * 1024 * 1024 # skip validation if > 64 MB + inputs_hex = [secrets.token_hex(sz) for _ in range(N)] + host_out, host_t = host_bench(inputs_hex, OUT_BYTES) + dev_out, dev_t, timing = device_bench(inputs_hex, OUT_BYTES) + match = (host_out == dev_out) if validate else (len(host_out) == len(dev_out)) speedup = host_t / dev_t if dev_t else float("inf") - print(f"{N:>8} {host_t*1e3:>10.1f} {dev_t*1e3:>10.1f} {speedup:>7.2f}x {match}") + total_str = f"{total_bytes/1e6:.1f}MB" if total_bytes >= 1e6 else f"{total_bytes/1e3:.0f}KB" + print(f"{sz:>10} {N:>9,} {total_str:>9} {host_t*1e3:>9.0f} {dev_t*1e3:>9.0f} " + f"{timing[2]:>10.2f} {speedup:>7.2f}x {match}") + if speedup > 1.0: + sweet.append((sz, N, total_bytes, host_t, dev_t, timing[2])) + + print() + if sweet: + print("=== GPU wins (speedup > 1.0×) ===") + for sz, N, tot, ht, dt, km in sweet: + print(f" input={sz:>8} × N={N:>7,} ({tot/1e6:>6.1f} MB):" + f" host {ht*1e3:>6.0f} ms → device {dt*1e3:>6.0f} ms ({ht/dt:.2f}×)") + print() + else: + print("=== GPU wins zone NOT REACHED ===") + print() + + print("=== honest breakdown (256 KB × 1,000 = 250 MB workload) ===") + inputs_hex = [secrets.token_hex(256 * 1024) for _ in range(1000)] + _, _, timing = device_bench(inputs_hex, OUT_BYTES) + ld, lc, kn, wr = timing + print(f" portal load : {ld:>8.1f} ms") + print(f" cuda init+upl : {lc:>8.1f} ms ← fixed cost per process spawn") + print(f" kernel : {kn:>8.2f} ms ← actual compute") + print(f" portal write : {wr:>8.1f} ms") + if kn > 0: + print(f" → kernel is {lc/kn:.0f}× faster than the launch/init phase") + print() + print("=== takeaway ===") + print(" - Per-spawn break-even is ~100 MB of bulk SHAKE work.") + print(" - Below that, host hashlib wins (no 200 ms init).") + print(" - Above that, kernel dominates & 3090 wins by 3–10×.") + print(" - For workloads naturally below the threshold:") + print(" (a) batch many fan-outs into one spawn, OR") + print(" (b) persistent daemon mode (skips cuda init entirely)") if __name__ == "__main__": diff --git a/examples/cuda-fanout/bench_daemon.py b/examples/cuda-fanout/bench_daemon.py new file mode 100644 index 0000000..42e5c34 --- /dev/null +++ b/examples/cuda-fanout/bench_daemon.py @@ -0,0 +1,107 @@ +"""bench_daemon.py — daemon-mode vs per-spawn comparison. + +Per-spawn pays ~200 ms CUDA context init every call. Daemon pays it +once at startup, then every subsequent call is just kernel + I/O. + +For workloads doing many small calls (the natural shape of a +go-gpu primitive in a search loop), daemon mode is the only mode +that wins.""" + +import hashlib +import os +import secrets +import subprocess +import sys +import time + +BIN = sys.argv[1] if len(sys.argv) > 1 else "./shake256-fanout" +OUT_BYTES = 32 + + +def write_portal(path, inputs_hex, out_bytes): + with open(path, "w") as fh: + fh.write("(cuda-shake-fanout\n") + fh.write(f" (output-bytes {out_bytes})\n") + fh.write(" (inputs\n") + for h in inputs_hex: + fh.write(f' "{h}"\n') + fh.write("))\n") + + +def main(): + print(f"binary: {BIN}\n") + N_INPUTS = 100 # per-call (small workload) + N_CALLS = 10 # number of fan-out calls + + # generate N_CALLS distinct input sets + portals = [] + for i in range(N_CALLS): + hex_in = [secrets.token_hex(32) for _ in range(N_INPUTS)] + in_p = f"/tmp/cf-d-in-{i}.portal" + write_portal(in_p, hex_in, OUT_BYTES) + portals.append((in_p, f"/tmp/cf-d-out-{i}.portal")) + + # ── per-spawn (control) ──────────────────────────────────────── + t0 = time.monotonic() + for in_p, out_p in portals: + subprocess.run([BIN, in_p, out_p], check=True) + spawn_t = time.monotonic() - t0 + + # ── daemon mode ──────────────────────────────────────────────── + t0 = time.monotonic() + proc = subprocess.Popen([BIN, "--daemon"], + stdin=subprocess.PIPE, stdout=subprocess.PIPE, text=True) + ready = proc.stdout.readline().strip() + assert ready == "ready", f"daemon didn't ready: {ready!r}" + init_t = time.monotonic() - t0 + + t0 = time.monotonic() + for in_p, out_p in portals: + proc.stdin.write(f"process {in_p} {out_p}\n") + proc.stdin.flush() + line = proc.stdout.readline().strip() + assert line.startswith("done"), f"daemon error: {line!r}" + daemon_calls_t = time.monotonic() - t0 + + proc.stdin.write("quit\n"); proc.stdin.flush() + proc.stdout.readline(); proc.wait() + daemon_total = init_t + daemon_calls_t + + # ── host hashlib (reference) ─────────────────────────────────── + inputs_hex = [secrets.token_hex(32) for _ in range(N_INPUTS)] + t0 = time.monotonic() + for _ in range(N_CALLS): + [hashlib.shake_256(bytes.fromhex(h)).hexdigest(OUT_BYTES) for h in inputs_hex] + host_t = time.monotonic() - t0 + + # cleanup + for in_p, out_p in portals: + try: os.unlink(in_p) + except: pass + try: os.unlink(out_p) + except: pass + + # ── report ───────────────────────────────────────────────────── + print(f"workload: {N_CALLS} calls × {N_INPUTS} inputs × 32 bytes each\n") + print(f" host hashlib loop : {host_t*1e3:>8.1f} ms total ({host_t*1e3/N_CALLS:.2f} ms/call)") + print(f" per-spawn fanout : {spawn_t*1e3:>8.1f} ms total ({spawn_t*1e3/N_CALLS:.2f} ms/call)") + print(f" daemon-mode init : {init_t*1e3:>8.1f} ms (one-time)") + print(f" daemon-mode calls : {daemon_calls_t*1e3:>8.1f} ms total " + f"({daemon_calls_t*1e3/N_CALLS:.2f} ms/call)") + print(f" daemon end-to-end : {daemon_total*1e3:>8.1f} ms (init + calls)") + print() + print(f" daemon/per-spawn : {daemon_total/spawn_t:.3f}× ({(1 - daemon_total/spawn_t)*100:+.1f}% wall time)") + print(f" daemon/host : {daemon_total/host_t:.2f}× " + f"({'GPU wins' if daemon_total < host_t else 'host wins'})") + print() + if N_CALLS >= 2: + avg_per_call = daemon_calls_t / N_CALLS + spawn_per_call = spawn_t / N_CALLS + print(f" per-call comparison (excluding daemon init):") + print(f" per-spawn: {spawn_per_call*1e3:.2f} ms each (cuda ctx init each)") + print(f" daemon : {avg_per_call*1e3:.2f} ms each (warm context)") + print(f" daemon is {spawn_per_call/avg_per_call:.1f}× faster per call") + + +if __name__ == "__main__": + main() diff --git a/examples/cuda-fanout/shake256-fanout.cu b/examples/cuda-fanout/shake256-fanout.cu index f28f9ee..43c66bd 100644 --- a/examples/cuda-fanout/shake256-fanout.cu +++ b/examples/cuda-fanout/shake256-fanout.cu @@ -208,20 +208,80 @@ static int parse_input_portal(const char *src, uint32_t *out_bytes, return count > 0; } +static int process_one(const char *in_path, const char *out_path); + int main(int argc, char **argv) { + /* daemon mode: stay alive across many calls, kill the 200 ms + * cuda-ctx-init overhead per spawn. + * + * Protocol (stdin, one command per line): + * process + * quit + * Stdout response (one line per command): + * done → success + * error → failure + * bye → quit acknowledged + */ + if (argc >= 2 && !strcmp(argv[1], "--daemon")) { + /* Touch the CUDA context once at startup so first request is fast. */ + int dev_count = 0; + cudaGetDeviceCount(&dev_count); + if (dev_count == 0) { + fprintf(stderr, "no CUDA devices\n"); return 1; + } + cudaFree(0); /* forces context init */ + fprintf(stdout, "ready\n"); fflush(stdout); + + char line[4096]; + while (fgets(line, sizeof(line), stdin)) { + char *nl = strchr(line, '\n'); + if (nl) *nl = 0; + if (!strcmp(line, "quit")) { + fprintf(stdout, "bye\n"); fflush(stdout); + break; + } + if (!strncmp(line, "process ", 8)) { + char *in_path = line + 8; + char *sep = strchr(in_path, ' '); + if (!sep) { + fprintf(stdout, "error bad-args\n"); fflush(stdout); + continue; + } + *sep = 0; + char *out_path = sep + 1; + if (process_one(in_path, out_path) == 0) { + fprintf(stdout, "done %s\n", out_path); + } else { + fprintf(stdout, "error process-failed\n"); + } + fflush(stdout); + } else if (line[0]) { + fprintf(stdout, "error unknown-command\n"); fflush(stdout); + } + } + return 0; + } + if (argc < 3) { - fprintf(stderr, "usage: %s \n", argv[0]); + fprintf(stderr, + "usage: %s \n" + " %s --daemon (read commands on stdin)\n", + argv[0], argv[0]); return 1; } + return process_one(argv[1], argv[2]); +} + +static int process_one(const char *in_path, const char *out_path) { struct timespec t0, t1, t2, t3, t4; clock_gettime(CLOCK_MONOTONIC, &t0); size_t portal_len; - char *portal = slurp(argv[1], &portal_len); + char *portal = slurp(in_path, &portal_len); uint32_t out_bytes = 32; char **hex_in = NULL; uint32_t n = 0; if (!parse_input_portal(portal, &out_bytes, &hex_in, &n)) { - fprintf(stderr, "no (inputs ...) found in %s\n", argv[1]); + fprintf(stderr, "no (inputs ...) found in %s\n", in_path); return 1; } @@ -268,8 +328,8 @@ int main(int argc, char **argv) { CUDA_CHECK(cudaMemcpy(h_out, d_out, (size_t)n * out_bytes, cudaMemcpyDeviceToHost)); /* write portal */ - FILE *f = fopen(argv[2], "w"); - if (!f) { perror(argv[2]); return 1; } + FILE *f = fopen(out_path, "w"); + if (!f) { perror(out_path); return 1; } fprintf(f, "(cuda-shake-fanout-result\n"); fprintf(f, " (n %u)\n", n); fprintf(f, " (output-bytes %u)\n", out_bytes); @@ -295,3 +355,4 @@ int main(int argc, char **argv) { free(hex_buf); free(portal); return 0; } +