engineering/lumbda
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

examples: cuda-fanout reference primitive for cross-tier GPU work

Browse source 
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.
This commit is contained in:
russell@unturf.com 2026-06-04 17:50:48 -04:00
parent ebcf1625c7
commit f11e310b24
No known key found for this signature in database
6 changed files with 667 additions and 0 deletions
Download patch file Download diff file Expand all files Collapse all files

39
examples/cuda-fanout/Makefile Normal file
View file

@ -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

128
examples/cuda-fanout/README.md Normal file
View file

@ -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.

68
examples/cuda-fanout/bench.py Normal file
View file

@ -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()

84
examples/cuda-fanout/lumbda-call.lsp Normal file
View file

@ -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

297
examples/cuda-fanout/shake256-fanout.cu Normal file
View file

@ -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 <input.portal> <output.portal>
*
* 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 <cuda_runtime.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <time.h>
#include <stdint.h>
#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 <input.portal> <output.portal>\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<<<blocks, threads>>>(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;
}

51
examples/cuda-fanout/test_roundtrip.py Normal file
View file

@ -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()