diff --git a/examples/cuda-fanout/README.md b/examples/cuda-fanout/README.md index d8743c8..e1daed5 100644 --- a/examples/cuda-fanout/README.md +++ b/examples/cuda-fanout/README.md @@ -197,6 +197,43 @@ A few cases: In those cases the per-spawn binary is sufficient and the daemon is overkill. +### Binary portal + daemon — GPU wins by 1.5–10× + +Adding a length-prefixed binary input/output format (`--binary` flag +and `process-bin` daemon command) eliminates the hex-string parse +that was eating 99% of wall time at large workloads. With **both +binary format and daemon mode**, every cell of the bench grid flips +to a GPU win: + +``` +in_sz N total host ms dev ms speedup +32 10,000 320 KB 5 1 5.45× +32 100,000 3.2 MB 47 5 9.95× +32 1,000,000 32 MB 470 47 10.11× +1024 100,000 102 MB 177 79 2.24× +16384 10,000 164 MB 231 124 1.86× +262144 1,000 262 MB 363 231 1.57× +``` + +Peak 10× at small-input × high-N. This is the natural shape of many +crypto protocols (commitments, Fiat-Shamir transforms, PoW search) +where each call hashes a small message but you make a lot of them. +The lumbda `(bend …)` primitive should default to binary format +under the hood; the S-expression layer is just the friendly client +contract. + +Binary format wire (in): +``` +u32 LE out_bytes +u32 LE n +for each i: u32 LE len_i | len_i bytes +``` + +Binary format wire (out): +``` +u32 LE n | u32 LE out_bytes | n × out_bytes +``` + ### Workloads that genuinely shine on this architecture SHAKE256 is too light. Workloads that actually win on GPU pour diff --git a/examples/cuda-fanout/bench_binary.py b/examples/cuda-fanout/bench_binary.py new file mode 100644 index 0000000..713b972 --- /dev/null +++ b/examples/cuda-fanout/bench_binary.py @@ -0,0 +1,124 @@ +"""bench_binary.py — binary portal format vs text portal. + +Same workloads as bench.py, but uses the new binary input/output +format that skips hex parsing entirely. This is where GPU should +actually start winning at heavy workloads.""" + +import hashlib +import os +import secrets +import struct +import subprocess +import sys +import time + +BIN = sys.argv[1] if len(sys.argv) > 1 else "./shake256-fanout" +OUT_BYTES = 32 + + +def write_binary_input(path, inputs_bytes, out_bytes): + """Format: u32 out_bytes | u32 n | (u32 len | len bytes) × n""" + with open(path, "wb") as fh: + fh.write(struct.pack("10} {'N':>9} {'total':>9} {'host ms':>9} " + f"{'dev (1shot)':>11} {'dev (daem)':>10} {'spd_daem':>9} match") + print("-" * 94) + sweet = [] + for sz, N in GRID: + inputs_bytes = [secrets.token_bytes(sz) for _ in range(N)] + host_out, host_t = host_bench(inputs_bytes, OUT_BYTES) + # daemon + dev_d_out, dev_d_t = device_bench(inputs_bytes, OUT_BYTES, use_daemon_proc=proc) + # one-shot + dev_1_out, dev_1_t = device_bench(inputs_bytes, OUT_BYTES) + match = host_out == dev_d_out == dev_1_out + speedup_d = host_t / dev_d_t if dev_d_t else float("inf") + total = sz * N + total_str = f"{total/1e6:.1f}MB" if total >= 1e6 else f"{total/1e3:.0f}KB" + print(f"{sz:>10} {N:>9,} {total_str:>9} {host_t*1e3:>9.0f} " + f"{dev_1_t*1e3:>11.0f} {dev_d_t*1e3:>10.0f} {speedup_d:>8.2f}x {match}") + if speedup_d > 1.0: + sweet.append((sz, N, total, host_t, dev_d_t)) + + proc.stdin.write("quit\n"); proc.stdin.flush() + proc.stdout.readline(); proc.wait() + + print() + if sweet: + print("=== GPU wins (daemon mode, binary format) ===") + for sz, N, tot, ht, dt 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}×)") + else: + print("=== GPU still loses everywhere ===") + print() + print("Columns: dev (1shot) includes 200 ms cuda init each call.") + print(" dev (daem) reuses warm daemon — the real production cost.") + + +if __name__ == "__main__": + main() diff --git a/examples/cuda-fanout/bend.lsp b/examples/cuda-fanout/bend.lsp new file mode 100644 index 0000000..1f7ae2a --- /dev/null +++ b/examples/cuda-fanout/bend.lsp @@ -0,0 +1,169 @@ +;;; bend.lsp — Lisp-smart GPU dispatch primitive. +;;; +;;; (bend expr) +;;; → inspects expr at runtime; routes to a registered GPU worker +;;; if the cost estimator says it's worth the trip, else evaluates +;;; locally in the original lexical scope. +;;; +;;; The decision is per-call, based on what's in the args, not just +;;; the op name. Tiny SHAKE on 5 inputs stays local; big SHAKE on +;;; 5,000,000 inputs bends to a GPU worker. +;;; +;;; Usage: +;;; +;;; (bend (cuda-shake-fanout small-list 32)) +;;; → local (cost below threshold) +;;; +;;; (bend (cuda-shake-fanout huge-list 32)) +;;; → GPU dispatch (cost above threshold) +;;; +;;; (bend! (cuda-shake-fanout x 32)) +;;; → force GPU, error if no worker available +;;; +;;; Per-tier integration: each lumbda tier registers `bend` in its +;;; primitive dispatch. Python tier uses subprocess.run + TCP socket +;;; module; C tier uses fork + portal; asm tier uses syscall fork + +;;; sock_stream. The Scheme wrapper below is the SHAPE; tier-specific +;;; details land in each tier's bend.{py,c,s}. + +;;; ── registry ────────────────────────────────────────────────── +;;; +;;; Whitelist of op names that *can* run on a GPU worker. Each entry +;;; carries (op-name . cost-estimator) where cost-estimator is a +;;; function that takes the literal arg list and returns an estimate +;;; in nanoseconds. Comparing estimate vs *bend-overhead-ns* decides. + +(define *bend-overhead-ns* 200000000) ; 200 ms — per-spawn cuda init + ; with daemon up, set to 100000 (100 µs RPC) + +(define *bend-host-throughput-bytes-per-ns* 0.5) ; ~500 MB/s host SHAKE256 + +(define *bend-gpu-ops* '()) ; alist of (op . cost-fn) + +(define (bend-register-op! name cost-fn) + "Add a GPU-capable op to the registry." + (set! *bend-gpu-ops* (cons (cons name cost-fn) *bend-gpu-ops*))) + +;;; Cost estimator for cuda-shake-fanout: (N inputs × avg-len) bytes +(define (bend-shake-cost args) + (let* ((inputs (car args)) + (out-bytes (car (cdr args))) + (n (length inputs)) + ;; Each input is a hex string; bytes = strlen / 2. + (avg-len (if (= n 0) 0 + (quotient (string-length (car inputs)) 2))) + (total-bytes (* n avg-len))) + ;; host_ns = total_bytes / 0.5 ns/byte + (quotient total-bytes + (let ((thr *bend-host-throughput-bytes-per-ns*)) + (if (= thr 0) 1 thr))))) + +(bend-register-op! 'cuda-shake-fanout bend-shake-cost) + +;;; ── worker connection ───────────────────────────────────────── + +(define *bend-worker-host* "localhost") +(define *bend-worker-port* 9091) + +(define (bend-worker-available?) + "Cheap probe — does a TCP connect to (host, port) succeed? + Each tier implements tcp-try-connect with whatever its socket + primitive is (Python socket, C connect(), asm syscall)." + (tcp-try-connect *bend-worker-host* *bend-worker-port*)) + +(define (bend-dispatch-to-gpu quoted-form) + "Serialize the quoted form as S-expression text, send to worker, + read back a result S-expression. Worker is responsible for matching + the op name against its local cuda binary." + (let ((sock (tcp-connect *bend-worker-host* *bend-worker-port*))) + (tcp-send sock (sexp->string quoted-form)) + (tcp-send sock "\n") + (let ((response (read-from-string (tcp-recv-until sock #\newline)))) + (tcp-close sock) + ;; response shape: (ok result) or (error reason) + (cond + ((eq? (car response) 'ok) (car (cdr response))) + ((eq? (car response) 'error) + (error "bend worker error:" (cdr response))) + (else (error "unexpected worker response:" response)))))) + +;;; ── core dispatcher ─────────────────────────────────────────── + +(define (bend-dispatch thunk quoted-form force-gpu?) + "Decide whether to evaluate locally or ship to a GPU worker. + thunk is a 0-arg lambda that, when called, evaluates the form in + its original lexical scope. quoted-form is the literal sexp, + used for cost estimation + GPU dispatch." + (let* ((head (if (pair? quoted-form) (car quoted-form) #f)) + (entry (assoc head *bend-gpu-ops*))) + (cond + ;; Force-mode: must go to GPU, error otherwise. + (force-gpu? + (if (and entry (bend-worker-available?)) + (bend-dispatch-to-gpu quoted-form) + (error "bend!: no GPU worker available for op" head))) + ;; Not a registered GPU op → local. + ((not entry) (thunk)) + (else + (let* ((cost-fn (cdr entry)) + (args (cdr quoted-form)) + (est-host-ns (cost-fn args))) + (cond + ;; Cost below overhead → local always wins. + ((< est-host-ns *bend-overhead-ns*) + (thunk)) + ;; Cost above overhead AND worker reachable → bend. + ((bend-worker-available?) + (bend-dispatch-to-gpu quoted-form)) + ;; Worker unreachable → fall back to local & note it. + (else + (display "bend: worker unreachable, running local: ") + (display head) (newline) + (thunk)))))))) + +;;; ── macros ──────────────────────────────────────────────────── + +(define-syntax bend + (syntax-rules () + ((_ expr) + (bend-dispatch (lambda () expr) 'expr #f)))) + +(define-syntax bend! + (syntax-rules () + ((_ expr) + (bend-dispatch (lambda () expr) 'expr #t)))) + +;;; ── helpers each tier must provide ──────────────────────────── +;;; +;;; The Scheme here is portable. These helpers are tier-specific: +;;; +;;; (tcp-try-connect host port) → #t / #f, no exception +;;; (tcp-connect host port) → socket handle +;;; (tcp-send sock str) +;;; (tcp-recv-until sock delim) → string +;;; (tcp-close sock) +;;; (sexp->string sexp) → S-expression as text +;;; +;;; Lumbda's fleet code (`ecdsa/lumbda/fleet/`) already implements all +;;; of these in Scheme — we should lift them out into a stdlib later. + +;;; ── demo ────────────────────────────────────────────────────── + +(define (demo) + ;; Small input — should stay local (below threshold). + (display "small input → ") + (let ((small (bend (cuda-shake-fanout '("00" "01" "deadbeef") 32)))) + (display (length small)) (display " hashes\n")) + ;; Larger input — should bend to GPU when worker is up. + ;; (Mocking with 10K identical inputs for cost-only test.) + (let ((huge (make-huge-input 10000))) + (display "huge input → ") + (let ((result (bend (cuda-shake-fanout huge 32)))) + (display (length result)) (display " hashes\n")))) + +(define (make-huge-input n) + (let loop ((i 0) (acc '())) + (if (= i n) acc + (loop (+ i 1) (cons "deadbeefcafebabe1234567890abcdef" acc))))) + +;; (demo) ; uncomment when tier-specific helpers + GPU worker are wired diff --git a/examples/cuda-fanout/gpu-worker.lsp b/examples/cuda-fanout/gpu-worker.lsp new file mode 100644 index 0000000..3f53049 --- /dev/null +++ b/examples/cuda-fanout/gpu-worker.lsp @@ -0,0 +1,158 @@ +;;; gpu-worker.lsp — TCP listener that dispatches bend-forms to local +;;; CUDA daemon binaries. +;;; +;;; Architecture: +;;; +;;; ┌─── client tier (any lumbda) ────┐ +;;; │ (bend (cuda-shake-fanout ...)) │ +;;; │ ↓ tcp-connect │ +;;; └─────────────┬───────────────────┘ +;;; │ S-expression text + newline +;;; ┌─────────────▼───────────────────┐ +;;; │ this gpu-worker on port 9091 │ +;;; │ - matches op name to local │ +;;; │ binary via capability list │ +;;; │ - writes input portal │ +;;; │ - sends "process in out" to │ +;;; │ the cuda-shake-fanout daemon │ +;;; │ - reads output portal │ +;;; │ - returns (ok result) text │ +;;; └─────────────────────────────────┘ +;;; +;;; Capability discovery: binaries advertised in *capabilities* list. +;;; Each entry: (op-name . daemon-handle). daemon-handle is a stdio +;;; pair to a long-lived --daemon-mode binary (CUDA context stays +;;; warm across all requests this worker handles). +;;; +;;; Run: +;;; +;;; lumbda gpu-worker.lsp ; default port 9091 +;;; LUMBDA_GPU_PORT=9001 lumbda gpu-worker.lsp + +(load "bend.lsp") ;; just for the registry; could be split + +(define *worker-port* 9091) +(define *capabilities* '()) ; alist (op-name . daemon-pipe) + +;;; ── daemon pool ─────────────────────────────────────────────── + +(define (start-daemon binary-path) + "Spawn `binary-path --daemon`, return (stdin-port . stdout-port). + stdout's first line must be 'ready' before we use it. + + Each tier implements (spawn-process-stdio path args) in its own + primitive (subprocess.Popen in Python, popen2 in C, syscall fork + + pipe in asm)." + (let* ((pair (spawn-process-stdio binary-path '("--daemon"))) + (ready (read-line (cdr pair)))) + (if (string=? ready "ready") + pair + (error "daemon failed to ready:" ready)))) + +(define (daemon-process daemon-pair in-portal out-portal) + "Send 'process in out' to a warm daemon, await 'done out' or 'error'." + (let ((in (car daemon-pair)) + (out (cdr daemon-pair))) + (display "process " in) (display in-portal in) + (display " " in) (display out-portal in) + (display "\n" in) + (flush-port in) + (let ((line (read-line out))) + (cond + ((string-prefix? "done " line) 'ok) + ((string-prefix? "error" line) (cons 'error line)) + (else (cons 'error (string-append "unknown response: " line))))))) + +(define (register-capability! op-name binary-path) + "Spawn a warm daemon for op-name, add to capability list." + (let ((daemon (start-daemon binary-path))) + (set! *capabilities* + (cons (cons op-name daemon) *capabilities*)) + (display "ready: ") (display op-name) + (display " ← ") (display binary-path) (newline))) + +;;; ── op handlers ─────────────────────────────────────────────── + +(define (handle-cuda-shake-fanout args) + "args = (inputs out-bytes). Write portal, dispatch to daemon, read + portal, return (ok result)." + (let* ((inputs (car args)) + (out-bytes (car (cdr args))) + (in-path (gensym-path "/tmp/bend-in" ".portal")) + (out-path (gensym-path "/tmp/bend-out" ".portal"))) + (write-shake-input-portal! in-path inputs out-bytes) + (let* ((daemon (cdr (assoc 'cuda-shake-fanout *capabilities*))) + (status (daemon-process daemon in-path out-path))) + (cond + ((eq? status 'ok) + (let ((result (read-shake-output-portal out-path))) + (delete-file in-path) + (delete-file out-path) + (list 'ok result))) + (else + (delete-file in-path) (delete-file out-path) + (list 'error (cdr status))))))) + +(define (write-shake-input-portal! path inputs out-bytes) + (let ((port (open-output-file path))) + (display "(cuda-shake-fanout\n" port) + (display " (output-bytes " port) (display out-bytes port) (display ")\n" port) + (display " (inputs\n" port) + (for-each (lambda (h) + (display " \"" port) (display h port) (display "\"\n" port)) + inputs) + (display "))\n" port) + (close-port port))) + +(define (read-shake-output-portal path) + ;; sim.lsp-style portal reader; just extract the hashes list. + (let ((sexp (read-from-string (file->string path)))) + ;; sexp shape: (cuda-shake-fanout-result ...) + ;; pull (hashes "..." "..." ...) section + (let loop ((children (cdr sexp))) + (cond + ((null? children) '()) + ((and (pair? (car children)) (eq? (car (car children)) 'hashes)) + (cdr (car children))) + (else (loop (cdr children))))))) + +;;; ── dispatch ────────────────────────────────────────────────── + +(define (handle-request sexp) + "Parse an incoming bend-form, look up its handler, run it. + Returns (ok ...) or (error reason)." + (cond + ((not (pair? sexp)) (list 'error "not a form")) + (else + (let* ((op (car sexp)) + (args (cdr sexp))) + (cond + ((eq? op 'cuda-shake-fanout) (handle-cuda-shake-fanout args)) + (else (list 'error (string-append "unknown op: " + (symbol->string op))))))))) + +(define (worker-loop sock) + "Accept loop. Each client sends one S-expression + newline; we reply + one S-expression + newline; close." + (let* ((client (tcp-accept sock)) + (line (tcp-recv-until client #\newline)) + (sexp (read-from-string line)) + (resp (handle-request sexp))) + (tcp-send client (sexp->string resp)) + (tcp-send client "\n") + (tcp-close client) + (worker-loop sock))) + +;;; ── entry ───────────────────────────────────────────────────── + +(define (main) + ;; Discover binaries — for v1, hard-coded paths. Future: read + ;; ~/.lumbda/gpu-capabilities or scan PATH. + (register-capability! 'cuda-shake-fanout + "/usr/local/bin/cuda-shake-fanout") + (let ((sock (tcp-listen *worker-port*))) + (display "gpu-worker listening on port ") + (display *worker-port*) (newline) + (worker-loop sock))) + +;; (main) ; uncomment to run; daemon binaries must exist on PATH diff --git a/examples/cuda-fanout/shake256-fanout.cu b/examples/cuda-fanout/shake256-fanout.cu index 43c66bd..a703ad3 100644 --- a/examples/cuda-fanout/shake256-fanout.cu +++ b/examples/cuda-fanout/shake256-fanout.cu @@ -209,6 +209,7 @@ static int parse_input_portal(const char *src, uint32_t *out_bytes, } static int process_one(const char *in_path, const char *out_path); +static int process_one_bin(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 @@ -240,8 +241,12 @@ int main(int argc, char **argv) { fprintf(stdout, "bye\n"); fflush(stdout); break; } - if (!strncmp(line, "process ", 8)) { - char *in_path = line + 8; + int is_bin = 0; + char *cmd_args = NULL; + if (!strncmp(line, "process ", 8)) { cmd_args = line + 8; is_bin = 0; } + else if (!strncmp(line, "process-bin ", 12)) { cmd_args = line + 12; is_bin = 1; } + if (cmd_args) { + char *in_path = cmd_args; char *sep = strchr(in_path, ' '); if (!sep) { fprintf(stdout, "error bad-args\n"); fflush(stdout); @@ -249,7 +254,9 @@ int main(int argc, char **argv) { } *sep = 0; char *out_path = sep + 1; - if (process_one(in_path, out_path) == 0) { + int r = is_bin ? process_one_bin(in_path, out_path) + : process_one(in_path, out_path); + if (r == 0) { fprintf(stdout, "done %s\n", out_path); } else { fprintf(stdout, "error process-failed\n"); @@ -262,16 +269,115 @@ int main(int argc, char **argv) { return 0; } + if (argc >= 4 && !strcmp(argv[1], "--binary")) { + return process_one_bin(argv[2], argv[3]); + } if (argc < 3) { fprintf(stderr, "usage: %s \n" + " %s --binary \n" " %s --daemon (read commands on stdin)\n", - argv[0], argv[0]); + argv[0], argv[0], argv[0]); return 1; } return process_one(argv[1], argv[2]); } +/* Binary format: + * in: u32 out_bytes | u32 n | (u32 len | len bytes) × n + * out: u32 n | u32 out_bytes | (out_bytes bytes) × n + * + * Eliminates the hex-string parse + emit overhead. At 262 MB workloads + * we measured 421 sec parsing 524 MB of hex text; binary skips that + * entirely. */ +static int process_one_bin(const char *in_path, const char *out_path) { + struct timespec t0, t1, t2, t3, t4; + clock_gettime(CLOCK_MONOTONIC, &t0); + + FILE *f = fopen(in_path, "rb"); + if (!f) { perror(in_path); return 1; } + fseek(f, 0, SEEK_END); long sz = ftell(f); fseek(f, 0, SEEK_SET); + uint8_t *buf = (uint8_t*)malloc(sz); + if (fread(buf, 1, sz, f) != (size_t)sz) { perror(in_path); return 1; } + fclose(f); + + uint32_t out_bytes, n; + memcpy(&out_bytes, buf, 4); + memcpy(&n, buf + 4, 4); + /* slice into offsets / lengths */ + uint32_t *offsets = (uint32_t*)malloc(n * sizeof(uint32_t)); + uint32_t *lengths = (uint32_t*)malloc(n * sizeof(uint32_t)); + size_t cur = 8; + size_t data_off = 0; + /* one pass: collect lengths, compute offsets in a separate flat array */ + size_t total = 0; + for (uint32_t i = 0; i < n; i++) { + if (cur + 4 > (size_t)sz) { + fprintf(stderr, "%s: truncated at input %u\n", in_path, i); return 1; + } + uint32_t L; memcpy(&L, buf + cur, 4); cur += 4; + if (cur + L > (size_t)sz) { + fprintf(stderr, "%s: bad length at input %u\n", in_path, i); return 1; + } + offsets[i] = (uint32_t)data_off; + lengths[i] = L; + data_off += L; + cur += L; + } + total = data_off; + /* flat array — just point at the buf bytes via offsets relative to start */ + uint8_t *flat = (uint8_t*)malloc(total ? total : 1); + cur = 8; + size_t fcur = 0; + for (uint32_t i = 0; i < n; i++) { + uint32_t L = lengths[i]; + cur += 4; + memcpy(flat + fcur, buf + cur, L); + cur += L; fcur += L; + } + free(buf); + clock_gettime(CLOCK_MONOTONIC, &t1); + + 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)); + + FILE *fo = fopen(out_path, "wb"); + if (!fo) { perror(out_path); return 1; } + fwrite(&n, 4, 1, fo); + fwrite(&out_bytes, 4, 1, fo); + fwrite(h_out, 1, (size_t)n * out_bytes, fo); + fclose(fo); + 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(stderr, " bin: load=%.1f launch=%.1f kernel=%.2f write=%.1f ms (n=%u, %.0f MB)\n", + ld_ms, launch_ms, kern_ms, write_ms, n, total / 1e6); + + cudaFree(d_in); cudaFree(d_off); cudaFree(d_len); cudaFree(d_out); + free(offsets); free(lengths); free(flat); free(h_out); + return 0; +} + static int process_one(const char *in_path, const char *out_path) { struct timespec t0, t1, t2, t3, t4; clock_gettime(CLOCK_MONOTONIC, &t0);