From f11e310b24f36f5245cc70b0bab2abf67199c8d9 Mon Sep 17 00:00:00 2001 From: "russell@unturf.com" Date: Thu, 4 Jun 2026 17:50:48 -0400 Subject: [PATCH] examples: cuda-fanout reference primitive for cross-tier GPU work MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Establishes the integration pattern for lumbda's future cuda primitive across Python / C / asm tiers without dragging the CUDA toolchain into lumbda's core build. Shape: leaf binary that every tier spawns via its existing process- spawn primitive & talks to through S-expression input + output portals. Asm tier inherits via fork + execve syscalls; no libcudart linkage; no DKMS dependency at lumbda build time. Files: shake256-fanout.cu self-contained CUDA SHAKE256 fan-out, Keccak permutation derived from FIPS 202 reference (tiny-sha3 lineage, CC0 → re-licensed AGPLv3) Makefile nvcc build + make test + make bench test_roundtrip.py validates output byte-identical to hashlib.shake_256 bench.py device vs host throughput at N = 1k / 10k / 100k lumbda-call.lsp reference Scheme wrapper showing the (cuda-shake-fanout inputs out-bytes) API shape lumbda's core would dispatch to per-tier README.md full integration story, wire contract, the three changes each tier needs (~20 LoC each), generalization path for other CUDA primitives Tested on 3090-ai (RTX 3090): make test → PASS — 4 / 4 hashes byte-identical to hashlib.shake_256 Honest bench (32-byte inputs): N host (Python hashlib) device (kernel launch dominated) 1,000 0.6 ms 188.1 ms 10,000 5.9 ms 195.7 ms 100,000 58.7 ms 313.8 ms Useful primitive when inputs are larger (KB+) or N reaches millions; honest about the launch-overhead break-even point. This is the reference, not the win — the win is locking the API shape so each tier registers under one stable name. Provenance: extracted as the generic pattern from ~/git/www.foxhop.net/ecdsa/cuda/sim_gpu.cu where on-device SHAKE delivered 2.6× memory compression for batched reversible-circuit simulation. Re-shipping the primitive back to the lumbda repo so the ecosystem inherits the work. --- examples/cuda-fanout/Makefile | 39 ++++ examples/cuda-fanout/README.md | 128 ++++++++++ examples/cuda-fanout/bench.py | 68 ++++++ examples/cuda-fanout/lumbda-call.lsp | 84 +++++++ examples/cuda-fanout/shake256-fanout.cu | 297 ++++++++++++++++++++++++ examples/cuda-fanout/test_roundtrip.py | 51 ++++ 6 files changed, 667 insertions(+) create mode 100644 examples/cuda-fanout/Makefile create mode 100644 examples/cuda-fanout/README.md create mode 100644 examples/cuda-fanout/bench.py create mode 100644 examples/cuda-fanout/lumbda-call.lsp create mode 100644 examples/cuda-fanout/shake256-fanout.cu create mode 100644 examples/cuda-fanout/test_roundtrip.py diff --git a/examples/cuda-fanout/Makefile b/examples/cuda-fanout/Makefile new file mode 100644 index 0000000..5e3bcb0 --- /dev/null +++ b/examples/cuda-fanout/Makefile @@ -0,0 +1,39 @@ +# shake256-fanout — Makefile. +# +# Builds the reference cuda fan-out worker. Each lumbda tier (Python, +# C, asm) spawns this binary via its existing process-spawn primitive +# & talks to it through input + output portal files. No tier links +# against libcudart directly — the CUDA toolchain dependency stays +# isolated to this directory. +# +# Build: +# make all — compile shake256-fanout (needs nvcc) +# make test — round-trip test (host + device produce identical +# SHAKE256 outputs over a small synthetic input set) +# make bench — bench device vs host throughput at N = 100k + +NVCC ?= nvcc +NVCCFLAGS ?= -O3 -arch=sm_86 -Xcompiler="-O3 -Wall -Wextra" +NVCC_PATH ?= $(shell which nvcc 2>/dev/null || echo /usr/local/cuda/bin/nvcc) + +all: shake256-fanout + +shake256-fanout: shake256-fanout.cu + $(NVCC) $(NVCCFLAGS) -o $@ $< + +# round-trip: feed (cuda-shake-fanout (output-bytes 32) (inputs ...)) +# through the binary, compare against an external SHAKE256 reference +# (Python hashlib.shake_256). +test: shake256-fanout test_roundtrip.py + @printf '(cuda-shake-fanout\n (output-bytes 32)\n (inputs\n' > /tmp/cf-in.portal + @printf ' "00"\n "01"\n "deadbeef"\n "%s"\n' \ + "$$(python3 -c 'print("ab" * 100)')" >> /tmp/cf-in.portal + @printf '))\n' >> /tmp/cf-in.portal + ./shake256-fanout /tmp/cf-in.portal /tmp/cf-out.portal + python3 test_roundtrip.py /tmp/cf-in.portal /tmp/cf-out.portal + +bench: shake256-fanout bench.py + python3 bench.py ./shake256-fanout + +clean: + rm -f shake256-fanout /tmp/cf-in.portal /tmp/cf-out.portal diff --git a/examples/cuda-fanout/README.md b/examples/cuda-fanout/README.md new file mode 100644 index 0000000..de88cd6 --- /dev/null +++ b/examples/cuda-fanout/README.md @@ -0,0 +1,128 @@ +# cuda-fanout — reference CUDA primitive for lumbda + +A reference implementation of "fan out a workload across CUDA cores" +as a primitive usable from any lumbda tier (Python, C, asm) without +linking the CUDA toolchain into lumbda's core build. + +License: **AGPLv3** (matches lumbda). + +## Why this shape + +Lumbda's strength is tier portability: the same Scheme source runs +byte-identically across Python / C / asm implementations. Direct +libcudart linkage breaks this — the asm tier has no libc, the C tier +gains a hard dependency, and Python's path-of-least-resistance +(cupy / pycuda) doesn't match the other tiers' wire format. + +Instead, this primitive is structured as a **leaf binary that every +tier spawns via its existing process-spawn primitive**: + +``` + ┌───────── lumbda Python tier ─────────┐ + │ (cuda-shake-fanout '("00" "01") 32) │ + │ ↓ subprocess.run │ + │ shake256-fanout in.portal out.portal │ + └──────────────────────────────────────┘ + + ┌───────── lumbda C tier ──────────────┐ + │ (cuda-shake-fanout ...) │ + │ ↓ fork + execve │ + │ shake256-fanout in.portal out.portal │ + └──────────────────────────────────────┘ + + ┌───────── lumbda asm tier ────────────┐ + │ (cuda-shake-fanout ...) │ + │ ↓ syscall fork + execve │ + │ shake256-fanout in.portal out.portal │ + └──────────────────────────────────────┘ +``` + +Every tier inherits the GPU work without touching its build +dependencies. The CUDA toolchain stays isolated to this directory. + +## Wire contract + +**Input portal** (caller writes): +``` +(cuda-shake-fanout + (output-bytes 32) + (inputs + "deadbeef" + "00ff00ff" + ...)) +``` + +**Output portal** (binary writes): +``` +(cuda-shake-fanout-result + (n 2) + (output-bytes 32) + (hashes + "abc..." + "def...") + (timing-ms (load 0.4) (launch 0.1) (kernel 12.7) (write 0.3))) +``` + +The wire format is plain S-expressions so any tier's portal reader +already handles it (see `c/portal.c`, `lumbda.py` portal helpers, +`asm/portal.s`). + +## What's in the box + +| file | role | +|------|------| +| `shake256-fanout.cu` | CUDA kernel + host driver. Self-contained, no external EC math, just SHAKE256 fan-out. | +| `Makefile` | nvcc build + `make test` + `make bench` | +| `test_roundtrip.py` | validates output matches `hashlib.shake_256` (host reference) | +| `bench.py` | timing comparison: Python `hashlib` host loop vs CUDA fan-out at N = 1k / 10k / 100k | +| `lumbda-call.lsp` | reference Scheme wrapper showing how lumbda would expose this as `(cuda-shake-fanout inputs out-bytes)` once each tier registers it | + +## How lumbda's core would integrate this + +Three things, each ~20 lines per tier: + +1. **Stable primitive name & dispatch entry**: `(cuda-shake-fanout + inputs out-bytes)` in each tier's primitive table. +2. **Host detection**: at startup, check whether the binary exists + at a known path (e.g. `~/.local/bin/cuda-shake-fanout` or + `LUMBDA_CUDA_FANOUT_BIN`); if not, the primitive raises a clean + `cuda-not-available` error. +3. **The Scheme wrapper** in `lumbda-call.lsp` (or equivalent inlined + per tier). + +No tier needs nvcc to build. No tier needs libcudart at runtime. The +GPU stays an opt-in capability per host. + +## Generalizing past SHAKE256 + +`shake256-fanout` is a deliberately narrow example because SHAKE256 is +small, well-known, and easy to validate against `hashlib.shake_256`. +The same pattern extends to any embarrassingly-parallel CUDA workload +the lumbda ecosystem wants — propose new primitives per-workload: + +- `cuda-shake-fanout` — N hashes (this binary) +- `cuda-keccakp-fanout` — N raw Keccak permutations (future) +- `cuda-poly1305-fanout` — N MACs (future) + +Each one ships as its own leaf binary; each one registers under its +own primitive name; each one is independently AGPLv3 alongside the +rest of the repo. + +## Build & smoke test + +```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 +``` + +## Provenance + +Keccak permutation derived from FIPS 202 reference. The compact +distillation pattern (Markku-Juhani O. Saarinen's tiny-sha3, CC0 / +public domain) was the starting point and is re-licensed here under +AGPLv3 to align with lumbda. The first practical use was in +`~/git/www.foxhop.net/ecdsa/cuda/sim_gpu.cu` where on-device SHAKE +gave a 2.6× memory compression for batched reversible-circuit +simulation (writeup in `ecdsa/cuda/README.md`). This contribution +extracts the generic primitive from that workload-specific use. diff --git a/examples/cuda-fanout/bench.py b/examples/cuda-fanout/bench.py new file mode 100644 index 0000000..45ee5af --- /dev/null +++ b/examples/cuda-fanout/bench.py @@ -0,0 +1,68 @@ +"""bench.py — N-input fan-out: device vs host hashlib timing. + +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.""" + +import hashlib +import os +import secrets +import subprocess +import sys +import time + +BIN = sys.argv[1] if len(sys.argv) > 1 else "./shake256-fanout" + + +def host_bench(inputs, out_bytes): + t0 = time.monotonic() + out = [hashlib.shake_256(bytes.fromhex(h)).hexdigest(out_bytes) for h in inputs] + return out, time.monotonic() - t0 + + +def write_portal(path, inputs, 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: + fh.write(f' "{h}"\n') + fh.write("))\n") + + +def read_hashes(path): + import re + with open(path) as fh: + text = fh.read() + m = re.search(r"\(hashes\b([^)]*)\)", text, re.DOTALL) + return re.findall(r'"([^"]*)"', m.group(1)) if m else [] + + +def device_bench(inputs, out_bytes): + in_p = "/tmp/cf-bench-in.portal" + out_p = "/tmp/cf-bench-out.portal" + write_portal(in_p, inputs, out_bytes) + t0 = time.monotonic() + subprocess.run([BIN, in_p, out_p], check=True) + wall = time.monotonic() - t0 + out = read_hashes(out_p) + os.unlink(in_p); os.unlink(out_p) + return out, wall + + +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 + 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}") + + +if __name__ == "__main__": + main() diff --git a/examples/cuda-fanout/lumbda-call.lsp b/examples/cuda-fanout/lumbda-call.lsp new file mode 100644 index 0000000..2377ddf --- /dev/null +++ b/examples/cuda-fanout/lumbda-call.lsp @@ -0,0 +1,84 @@ +;;; lumbda-call.lsp — calling shake256-fanout from any lumbda tier. +;;; +;;; This is the API SHAPE lumbda's core would expose as a primitive. +;;; Each tier already has spawn-process + portal I/O primitives; the +;;; cuda fan-out primitive is a thin Scheme wrapper that composes them. +;;; +;;; Once locked, each tier (Python / C / asm) registers the same name +;;; in its primitive dispatch table. Asm tier uses fork + execve + +;;; portal-write syscalls — no libcudart linkage needed. Python & +;;; C tiers can use the same shape via subprocess / fork respectively, +;;; or graduate to a true FFI later. +;;; +;;; Usage: +;;; +;;; (cuda-shake-fanout +;;; '("00" "01" "deadbeef") ; list of hex-string inputs +;;; 32) ; output bytes per input +;;; ⇒ '("0xhash0" "0xhash1" "0xhash2") + +(define *cuda-shake-fanout-binary* + ;; Override per host; reference points at the example dir. + "/home/fox/git/lumbda/examples/cuda-fanout/shake256-fanout") + +(define (with-temp-files names body) + "Allocate gensym'd temp file paths, invoke body on them, cleanup + regardless of exit path. Each tier should implement gensym + + delete-file in its stdlib." + (let* ((paths (map (lambda (n) (string-append "/tmp/cf-" n "-" + (number->string (current-time-ms)))) + names)) + (result (apply body paths))) + (for-each delete-file paths) + result)) + +(define (write-input-portal! path inputs output-bytes) + (let ((port (open-output-file path))) + (display "(cuda-shake-fanout\n" port) + (display " (output-bytes " port) (display output-bytes port) (display ")\n" port) + (display " (inputs\n" port) + (for-each (lambda (hex) + (display " \"" port) (display hex port) (display "\"\n" port)) + inputs) + (display "))\n" port) + (close-port port))) + +(define (read-hashes-portal path) + ;; Use sim.lsp-style portal reader if available; fall back to text scan. + (let* ((src (file->string path)) + ;; Crude hex-string extractor — production code should use a real + ;; parser. The portal format is small & predictable so this works + ;; for the reference example. + (after-hashes (substring src + (+ (string-search "(hashes" src) 7) + (string-length src)))) + (let loop ((rest after-hashes) (acc '())) + (let ((open (string-search "\"" rest))) + (if (not open) (reverse acc) + (let* ((tail (substring rest (+ open 1) (string-length rest))) + (close (string-search "\"" tail))) + (if (not close) (reverse acc) + (loop (substring tail (+ close 1) (string-length tail)) + (cons (substring tail 0 close) acc))))))))) + +(define (cuda-shake-fanout inputs output-bytes) + "Hash each input on its own CUDA thread via the reference fan-out + binary. Inputs are hex strings; output is a list of hex hash strings." + (with-temp-files '("in.portal" "out.portal") + (lambda (in-path out-path) + (write-input-portal! in-path inputs output-bytes) + ;; spawn-process is each tier's process-spawn primitive name — + ;; substitute when lumbda's core dispatches under a stable name. + (spawn-process *cuda-shake-fanout-binary* (list in-path out-path)) + (read-hashes-portal out-path)))) + +;;; ── demo ────────────────────────────────────────────────────── + +(define (demo) + (let ((hashes (cuda-shake-fanout '("00" "01" "deadbeef") 32))) + (display "got ") (display (length hashes)) (display " hashes:\n") + (for-each (lambda (h) + (display " ") (display h) (newline)) + hashes))) + +;; (demo) ;; uncomment to run; needs spawn-process + file primitives diff --git a/examples/cuda-fanout/shake256-fanout.cu b/examples/cuda-fanout/shake256-fanout.cu new file mode 100644 index 0000000..f28f9ee --- /dev/null +++ b/examples/cuda-fanout/shake256-fanout.cu @@ -0,0 +1,297 @@ +/* shake256-fanout.cu — reference cuda fan-out worker for lumbda. + * + * Reads an S-expression input portal listing N inputs, hashes each + * with SHAKE256 on its own GPU thread, writes an output portal with + * N hashes. The "fan-out primitive" lumbda would expose via its + * spawn-process primitive across Python / C / asm tiers — every tier + * inherits this binary without dragging the CUDA toolchain into + * lumbda's core. + * + * Wire format (input portal — written by the caller tier): + * + * (cuda-shake-fanout + * (output-bytes 32) + * (inputs + * "hex-bytes-1" + * "hex-bytes-2" + * ...)) + * + * Wire format (output portal — written by this binary): + * + * (cuda-shake-fanout-result + * (n 17) + * (output-bytes 32) + * (hashes + * "hex-output-1" + * "hex-output-2" + * ...) + * (timing-ms (load 0.4) (launch 0.1) (kernel 12.7) (write 0.3))) + * + * Usage: shake256-fanout + * + * License: AGPLv3. Keccak permutation derived from FIPS 202 reference; + * the tiny-sha3 distillation (Markku-Juhani O. Saarinen, CC0 / public + * domain) was used as the starting point and re-licensed here under + * AGPLv3 to align with lumbda's own license. + */ + +#include +#include +#include +#include +#include +#include + +#define CUDA_CHECK(stmt) do { \ + cudaError_t err = (stmt); \ + if (err != cudaSuccess) { \ + fprintf(stderr, "CUDA error %s at %s:%d: %s\n", \ + #stmt, __FILE__, __LINE__, cudaGetErrorString(err)); \ + exit(1); \ + } \ +} while (0) + +#define KECCAK_RATE 136 /* SHAKE256 rate (1600 - 2*256) / 8 */ + +/* ── Keccak permutation, __host__ __device__ — runs on either ── */ + +__host__ __device__ static const uint64_t RC[24] = { + 0x0000000000000001ULL, 0x0000000000008082ULL, + 0x800000000000808AULL, 0x8000000080008000ULL, + 0x000000000000808BULL, 0x0000000080000001ULL, + 0x8000000080008081ULL, 0x8000000000008009ULL, + 0x000000000000008AULL, 0x0000000000000088ULL, + 0x0000000080008009ULL, 0x000000008000000AULL, + 0x000000008000808BULL, 0x800000000000008BULL, + 0x8000000000008089ULL, 0x8000000000008003ULL, + 0x8000000000008002ULL, 0x8000000000000080ULL, + 0x000000000000800AULL, 0x800000008000000AULL, + 0x8000000080008081ULL, 0x8000000000008080ULL, + 0x0000000080000001ULL, 0x8000000080008008ULL, +}; +__host__ __device__ static const int R[24] = { + 1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14, + 27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44, +}; +__host__ __device__ static const int PI[24] = { + 10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4, + 15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1, +}; + +__host__ __device__ static inline uint64_t rotl64(uint64_t x, int n) { + return (x << n) | (x >> (64 - n)); +} + +__host__ __device__ static void keccakf(uint64_t s[25]) { + uint64_t t, bc[5]; + for (int r = 0; r < 24; r++) { + for (int i = 0; i < 5; i++) + bc[i] = s[i] ^ s[i + 5] ^ s[i + 10] ^ s[i + 15] ^ s[i + 20]; + for (int i = 0; i < 5; i++) { + t = bc[(i + 4) % 5] ^ rotl64(bc[(i + 1) % 5], 1); + for (int j = 0; j < 25; j += 5) s[j + i] ^= t; + } + t = s[1]; + for (int i = 0; i < 24; i++) { + int j = PI[i]; + bc[0] = s[j]; + s[j] = rotl64(t, R[i]); + t = bc[0]; + } + for (int j = 0; j < 25; j += 5) { + for (int i = 0; i < 5; i++) bc[i] = s[j + i]; + for (int i = 0; i < 5; i++) + s[j + i] ^= (~bc[(i + 1) % 5]) & bc[(i + 2) % 5]; + } + s[0] ^= RC[r]; + } +} + +/* per-thread SHAKE256: absorb input, finalize, squeeze out_len bytes. */ +__device__ static void shake256_oneshot(const uint8_t *in, uint32_t in_len, + uint8_t *out, uint32_t out_len) +{ + uint8_t state[200] = {0}; + uint32_t cur = 0; + for (uint32_t i = 0; i < in_len; i++) { + state[cur++] ^= in[i]; + if (cur == KECCAK_RATE) { keccakf((uint64_t*)state); cur = 0; } + } + /* SHAKE pad */ + state[cur] ^= 0x1F; + state[KECCAK_RATE - 1] ^= 0x80; + keccakf((uint64_t*)state); + cur = 0; + for (uint32_t i = 0; i < out_len; i++) { + if (cur == KECCAK_RATE) { keccakf((uint64_t*)state); cur = 0; } + out[i] = state[cur++]; + } +} + +/* kernel: one thread per input. inputs are stored back-to-back with + * a per-input offset[] + len[] index. outputs are flat output_bytes × + * N. block = 64 threads (~enough warps for typical fan-out widths). */ +__global__ void shake_fanout_kernel(const uint8_t *inputs, + const uint32_t *offsets, + const uint32_t *lengths, + uint8_t *outputs, + uint32_t output_bytes, + uint32_t n) +{ + uint32_t t = blockIdx.x * blockDim.x + threadIdx.x; + if (t >= n) return; + shake256_oneshot(inputs + offsets[t], lengths[t], + outputs + (size_t)t * output_bytes, output_bytes); +} + +/* ── host: portal I/O + driver ────────────────────────────────────── */ + +static char *slurp(const char *path, size_t *len_out) { + FILE *f = fopen(path, "rb"); + if (!f) { perror(path); exit(1); } + fseek(f, 0, SEEK_END); long n = ftell(f); fseek(f, 0, SEEK_SET); + char *buf = (char*)malloc(n + 1); + if (fread(buf, 1, n, f) != (size_t)n) { perror(path); exit(1); } + buf[n] = 0; fclose(f); + if (len_out) *len_out = n; + return buf; +} + +static int hex_to_bytes(const char *hex, size_t hex_len, uint8_t *out) { + if (hex_len % 2 != 0) return 0; + for (size_t i = 0; i < hex_len; i += 2) { + unsigned v; + if (sscanf(hex + i, "%2x", &v) != 1) return 0; + out[i / 2] = (uint8_t)v; + } + return 1; +} + +static void bytes_to_hex(const uint8_t *bytes, size_t n, char *out) { + static const char *H = "0123456789abcdef"; + for (size_t i = 0; i < n; i++) { + out[2 * i] = H[bytes[i] >> 4]; + out[2 * i + 1] = H[bytes[i] & 0xF]; + } + out[2 * n] = 0; +} + +/* parse: extract `(output-bytes N)` & all "..."-quoted strings inside + * `(inputs ...)`. Tolerant of whitespace; assumes inputs are simple + * hex strings. */ +static int parse_input_portal(const char *src, uint32_t *out_bytes, + char ***hex_inputs, uint32_t *n_inputs) +{ + *out_bytes = 32; + const char *p = strstr(src, "(output-bytes"); + if (p) sscanf(p, "(output-bytes %u)", out_bytes); + p = strstr(src, "(inputs"); + if (!p) return 0; + char **list = NULL; + uint32_t count = 0, cap = 0; + p += strlen("(inputs"); + while ((p = strchr(p, '"')) != NULL) { + const char *end = strchr(p + 1, '"'); + if (!end) break; + if (count == cap) { + cap = cap ? cap * 2 : 16; + list = (char**)realloc(list, cap * sizeof(char*)); + } + size_t L = end - (p + 1); + list[count] = (char*)malloc(L + 1); + memcpy(list[count], p + 1, L); list[count][L] = 0; + count++; + p = end + 1; + } + *hex_inputs = list; + *n_inputs = count; + return count > 0; +} + +int main(int argc, char **argv) { + if (argc < 3) { + fprintf(stderr, "usage: %s \n", argv[0]); + return 1; + } + struct timespec t0, t1, t2, t3, t4; + clock_gettime(CLOCK_MONOTONIC, &t0); + + size_t portal_len; + char *portal = slurp(argv[1], &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]); + return 1; + } + + /* lay out flat input bytes with offsets+lengths */ + uint32_t *offsets = (uint32_t*)malloc(n * sizeof(uint32_t)); + uint32_t *lengths = (uint32_t*)malloc(n * sizeof(uint32_t)); + size_t total = 0; + for (uint32_t i = 0; i < n; i++) { + size_t hex_len = strlen(hex_in[i]); + if (hex_len % 2 != 0) { + fprintf(stderr, "input %u has odd hex length\n", i); return 1; + } + offsets[i] = total; + lengths[i] = hex_len / 2; + total += hex_len / 2; + } + uint8_t *flat = (uint8_t*)malloc(total ? total : 1); + for (uint32_t i = 0; i < n; i++) { + if (!hex_to_bytes(hex_in[i], strlen(hex_in[i]), flat + offsets[i])) { + fprintf(stderr, "bad hex at input %u\n", i); return 1; + } + } + clock_gettime(CLOCK_MONOTONIC, &t1); + + /* upload */ + uint8_t *d_in; uint32_t *d_off, *d_len; uint8_t *d_out; + CUDA_CHECK(cudaMalloc(&d_in, total ? total : 1)); + CUDA_CHECK(cudaMalloc(&d_off, n * sizeof(uint32_t))); + CUDA_CHECK(cudaMalloc(&d_len, n * sizeof(uint32_t))); + CUDA_CHECK(cudaMalloc(&d_out, (size_t)n * out_bytes)); + CUDA_CHECK(cudaMemcpy(d_in, flat, total ? total : 1, cudaMemcpyHostToDevice)); + CUDA_CHECK(cudaMemcpy(d_off, offsets, n * sizeof(uint32_t), cudaMemcpyHostToDevice)); + CUDA_CHECK(cudaMemcpy(d_len, lengths, n * sizeof(uint32_t), cudaMemcpyHostToDevice)); + clock_gettime(CLOCK_MONOTONIC, &t2); + + int threads = 64; + int blocks = (n + threads - 1) / threads; + shake_fanout_kernel<<>>(d_in, d_off, d_len, d_out, out_bytes, n); + CUDA_CHECK(cudaGetLastError()); + CUDA_CHECK(cudaDeviceSynchronize()); + clock_gettime(CLOCK_MONOTONIC, &t3); + + uint8_t *h_out = (uint8_t*)malloc((size_t)n * out_bytes); + 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; } + fprintf(f, "(cuda-shake-fanout-result\n"); + fprintf(f, " (n %u)\n", n); + fprintf(f, " (output-bytes %u)\n", out_bytes); + fprintf(f, " (hashes\n"); + char *hex_buf = (char*)malloc(out_bytes * 2 + 1); + for (uint32_t i = 0; i < n; i++) { + bytes_to_hex(h_out + (size_t)i * out_bytes, out_bytes, hex_buf); + fprintf(f, " \"%s\"\n", hex_buf); + } + fprintf(f, " )\n"); + clock_gettime(CLOCK_MONOTONIC, &t4); + double ld_ms = (t1.tv_sec - t0.tv_sec) * 1e3 + (t1.tv_nsec - t0.tv_nsec) / 1e6; + double launch_ms = (t2.tv_sec - t1.tv_sec) * 1e3 + (t2.tv_nsec - t1.tv_nsec) / 1e6; + double kern_ms = (t3.tv_sec - t2.tv_sec) * 1e3 + (t3.tv_nsec - t2.tv_nsec) / 1e6; + double write_ms = (t4.tv_sec - t3.tv_sec) * 1e3 + (t4.tv_nsec - t3.tv_nsec) / 1e6; + fprintf(f, " (timing-ms (load %.3f) (launch %.3f) (kernel %.3f) (write %.3f)))\n", + ld_ms, launch_ms, kern_ms, write_ms); + fclose(f); + + cudaFree(d_in); cudaFree(d_off); cudaFree(d_len); cudaFree(d_out); + for (uint32_t i = 0; i < n; i++) free(hex_in[i]); + free(hex_in); free(offsets); free(lengths); free(flat); free(h_out); + free(hex_buf); free(portal); + return 0; +} diff --git a/examples/cuda-fanout/test_roundtrip.py b/examples/cuda-fanout/test_roundtrip.py new file mode 100644 index 0000000..4f01fba --- /dev/null +++ b/examples/cuda-fanout/test_roundtrip.py @@ -0,0 +1,51 @@ +"""test_roundtrip.py — verify shake256-fanout matches hashlib.shake_256. + +Reads the same input portal the cuda binary consumed, hashes each +input with Python's reference SHAKE256, compares against the binary's +output portal. If every hash matches, the cuda port is byte-identical +to upstream SHAKE256.""" + +import hashlib +import re +import sys + + +def extract_strings(text, section): + """Pull "..."-quoted strings inside (section ...).""" + m = re.search(r"\(" + section + r"\b([^)]*(?:\)[^)]+)*)", text, re.DOTALL) + if not m: + return [] + return re.findall(r'"([^"]*)"', m.group(1)) + + +def main(): + in_path, out_path = sys.argv[1], sys.argv[2] + with open(in_path) as fh: + in_text = fh.read() + with open(out_path) as fh: + out_text = fh.read() + + m = re.search(r"\(output-bytes\s+(\d+)\)", in_text) + out_bytes = int(m.group(1)) if m else 32 + + inputs = extract_strings(in_text, "inputs") + hashes = extract_strings(out_text, "hashes") + + if len(inputs) != len(hashes): + print(f"COUNT MISMATCH: {len(inputs)} inputs vs {len(hashes)} hashes") + sys.exit(1) + + n_ok = 0 + for i, (hex_in, got_hex) in enumerate(zip(inputs, hashes)): + expected = hashlib.shake_256(bytes.fromhex(hex_in)).hexdigest(out_bytes) + if expected != got_hex: + print(f"FAIL #{i}: input={hex_in[:20]}...") + print(f" expected {expected}") + print(f" got {got_hex}") + sys.exit(1) + n_ok += 1 + print(f"PASS — {n_ok} / {n_ok} hashes byte-identical to hashlib.shake_256") + + +if __name__ == "__main__": + main()