From 05fa922d428b45544688fe645b343e4106220f3b Mon Sep 17 00:00:00 2001 From: "russell@unturf.com" Date: Fri, 5 Jun 2026 21:50:05 -0400 Subject: [PATCH] =?UTF-8?q?bend=20form=20G=20=E2=80=94=20cuda-radix-sort?= =?UTF-8?q?=20lands=20live=20(CUB=20DeviceRadixSort=20u64)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Thin CUDA binary wrapping cub::DeviceRadixSort::SortKeys on a 64-bit key stream. One op wired day-1 (0x01 sort-u64-asc); 0x02/0x03/0x04 slots reserved (desc, u32, key-value) for future builds. Wire stays distinct from existing forms: request: BSRT | u32 op_id | u32 n | u64[n] response: BSRR | u32 status | u32 n | u64[n] sorted asc Validated on 3090-ai.foxhop.net byte-identical to Python sorted() at n ∈ {32, 1k, 100k, 1M, 10M}. Bench at sustained throughput: n kernel_ms Gkeys/s 100,000 0.142 0.706 1,000,000 0.265 3.767 10,000,000 1.817 5.504 ~4x over the published Titan baseline (1.4 Gkeys/s) at saturation, matching CUB's expected Ampere scaling. gpu-worker.lsp learns handle-binary-sort + BSRT magic dispatch + maybe-register-daemon! for cuda-radix-sort (overridable via RADIX_SORT_WORKER env). Both 3090-ai (:9091) & ai (:9092) workers restarted; both log `ready cuda-radix-sort <- ./radix-sort`. 4090 (ai.foxhop.net) standalone --binary run OOMs on cudaMalloc when all four daemons are co-resident (secp256k1 daemon parks ~24 GiB on startup, leaving 47 MiB free). Pre-existing capacity constraint of the ai host, not a form-G defect; tracked in form-G-progress.md. CATALOG.md & www/bend.html live-forms table updated with measured 3090 numbers; Wave 1 surveyed row for G marked as promoted. --- examples/cuda-fanout/CATALOG.md | 1 + examples/cuda-fanout/Makefile | 14 +- examples/cuda-fanout/gpu-worker.lsp | 44 +++ examples/cuda-fanout/plans/form-G-progress.md | 52 ++++ examples/cuda-fanout/radix-sort.cu | 252 ++++++++++++++++++ .../test_radix_sort_known_answers.py | 184 +++++++++++++ www/bend.html | 8 +- 7 files changed, 553 insertions(+), 2 deletions(-) create mode 100644 examples/cuda-fanout/plans/form-G-progress.md create mode 100644 examples/cuda-fanout/radix-sort.cu create mode 100644 examples/cuda-fanout/test_radix_sort_known_answers.py diff --git a/examples/cuda-fanout/CATALOG.md b/examples/cuda-fanout/CATALOG.md index 6369801..57c33c1 100644 --- a/examples/cuda-fanout/CATALOG.md +++ b/examples/cuda-fanout/CATALOG.md @@ -24,6 +24,7 @@ form-status column says `planned` until numbers exist. | `cuda-bignum-cgbn` | `cgbn-batch-worker` | 3090 | 1.28 Gops/s kernel mod-mul at n=1M (256-bit, ~256x GMP CPU); 9 ops total | `BCGB` binary: op_id + bitwidth + n + modulus + a + b | | `cuda-secp256k1-batched-mul`| `secp256k1-batch-mul` | 3090 | v1 7.86 Mkeys/s @ n=1M; v3 (windowed-G w=4) **13.83 Mkeys/s @ n=1M** (1.76x v1; ~309x coincurve); Day-4 v4 (v3 ladder + Montgomery batch inv) regressed -12% vs v3 due to v2 serial-per-block walks; warp-scan Phase B/D refactor needed before v4 wins; daemon default flipped to `--window-w 4` | `BSCP` binary: scalars + base-point | | `cuda-sim-axis-flip` | `ecdsa/cuda/demo_axis` | 3090 | 217 Mops/s @ K=32 M=4 (per-candidate parallel); 23.7x over per-shot N=4 at same M; LOSES to per-shot by 14% at full N=128 saturation | `(cuda-sim-axis (variant-paths …) n-shots)` | +| `cuda-radix-sort` | `radix-sort` | 3090 | **5.50 Gkeys/s @ n=10M** (kernel 1.82 ms); 3.77 Gkeys/s @ n=1M; CUB DeviceRadixSort u64 ascending; ~4x over published Titan baseline (1.4 Gkeys/s); byte-identity vs Python sorted() at n∈{32, 1k, 100k, 1M, 10M} | `BSRT` binary: op_id + n + u64[n] → `BSRR` sorted u64[n] | `cuda-sim-ops-bin` numbers are humbling on purpose: kickmix is conditional-op heavy & branch-divergent. This catalog exists so diff --git a/examples/cuda-fanout/Makefile b/examples/cuda-fanout/Makefile index 9b361a0..27e3644 100644 --- a/examples/cuda-fanout/Makefile +++ b/examples/cuda-fanout/Makefile @@ -16,11 +16,23 @@ 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 cgbn-batch-worker secp256k1-batch-mul +all: shake256-fanout cgbn-batch-worker secp256k1-batch-mul radix-sort shake256-fanout: shake256-fanout.cu $(NVCC) $(NVCCFLAGS) -o $@ $< +# bend form G — batched u64 radix sort via NVIDIA CUB +# (header-only, ships with CUDA Toolkit so no extra deps). +# Wire: BSRT (sort-u64-asc) request → BSRR sorted-keys response. +radix-sort: radix-sort.cu + $(NVCC) $(NVCCFLAGS) --expt-relaxed-constexpr -o $@ $< + +radix-sort-test: radix-sort test_radix_sort_known_answers.py + python3 -u test_radix_sort_known_answers.py ./radix-sort + +radix-sort-bench: radix-sort test_radix_sort_known_answers.py + python3 -u test_radix_sort_known_answers.py ./radix-sort --bench + # bend form A — secp256k1 batched scalar*G via VanitySearch-Bitcrack # vendored GPUMath.h (AGPL-3.0). Per-thread Jacobian double-and-add, # per-thread inversion to affine. See vendor/vanity-search-bitcrack/NOTICE. diff --git a/examples/cuda-fanout/gpu-worker.lsp b/examples/cuda-fanout/gpu-worker.lsp index 218b2ff..bebba04 100644 --- a/examples/cuda-fanout/gpu-worker.lsp +++ b/examples/cuda-fanout/gpu-worker.lsp @@ -49,6 +49,13 @@ (or (get-environment-variable "SECP256K1_BATCH_WORKER") "./secp256k1-batch-mul")) +;; radix-sort — bend form G. Batched u64 ascending sort via NVIDIA CUB +;; DeviceRadixSort (header-only, ships with CUDA Toolkit). Binary wire +;; uses BSRT request / BSRR response magic, distinct from BSHK/BCGB/BSCP. +(define *binary-radix-sort* + (or (get-environment-variable "RADIX_SORT_WORKER") + "./radix-sort")) + (define (parse-port-arg args) (let loop ((rest args)) (cond @@ -357,6 +364,38 @@ (if (file-exists? out-path) (delete-file out-path)) (wire-send-raw client (string-append "BERR" (cdr status))))))))) +;; Binary wire for bend form G (cuda-radix-sort). +;; Payload begins with "BSRT"; pass entire blob through to the daemon, +;; which expects the same magic + header. Response begins with "BSRR" +;; on success or "BERR" on failure. +(define (handle-binary-sort client payload) + (let* ((daemon (cdr (assoc 'cuda-radix-sort *daemons*))) + (in-path (gensym-path "/tmp/bend-sort-in" ".bin")) + (out-path (gensym-path "/tmp/bend-sort-out" ".bin")) + (t-start (current-time-ms))) + (write-binary-file in-path payload) + (display ";;; bend RECV cuda-radix-sort bytes=") + (display (string-length payload)) + (display " t-ms=") (display t-start) (newline) + (let ((status (daemon-process daemon in-path out-path #t))) + (let ((wall-ms (- (current-time-ms) t-start))) + (cond + ((eq? status 'ok) + (let ((result-blob (read-binary-file out-path))) + (if (file-exists? in-path) (delete-file in-path)) + (if (file-exists? out-path) (delete-file out-path)) + (display ";;; bend DONE cuda-radix-sort") + (display " wall-ms=") (display wall-ms) + (display " out-bytes=") (display (string-length result-blob)) + (newline) + (wire-send-raw client result-blob))) + (else + (display ";;; bend FAIL cuda-radix-sort wall-ms=") + (display wall-ms) (newline) + (if (file-exists? in-path) (delete-file in-path)) + (if (file-exists? out-path) (delete-file out-path)) + (wire-send-raw client (string-append "BERR" (cdr status))))))))) + ;; Accept one client, handle one request, close. Returns #t to keep ;; serving, #f when the server should stop. (define (handle-one server) @@ -379,6 +418,10 @@ (string=? (substring payload 0 4) "BSCP")) (handle-binary-secp client payload) (tcp-close client) #t) + ((and (>= (string-length payload) 4) + (string=? (substring payload 0 4) "BSRT")) + (handle-binary-sort client payload) + (tcp-close client) #t) (else (let* ((req (read-from-string payload)) (resp (handle-request req))) @@ -420,6 +463,7 @@ (maybe-register-daemon! 'cuda-secp256k1-batched-mul *binary-secp256k1-batch* *secp-daemon-extra-args*) + (maybe-register-daemon! 'cuda-radix-sort *binary-radix-sort*) (let ((server (tcp-listen port))) (cond ((eq? server #f) diff --git a/examples/cuda-fanout/plans/form-G-progress.md b/examples/cuda-fanout/plans/form-G-progress.md new file mode 100644 index 0000000..c87b8bb --- /dev/null +++ b/examples/cuda-fanout/plans/form-G-progress.md @@ -0,0 +1,52 @@ +# form G — cuda-radix-sort — build progress + +## 2026-06-05 — landed live + +Thin wrapper around `cub::DeviceRadixSort::SortKeys` for u64 ascending +sort. Day-1 wire freezes one op (0x01 sort-u64-asc); reserves op slots +0x02 (desc), 0x03 (u32), 0x04 (key-value) in our dispatcher switch +without wiring them. + +### Wire + +- Request: `BSRT` (4 B) + u32 op_id + u32 n + u64[n] +- Response: `BSRR` (4 B) + u32 status + u32 n + u64[n] sorted ascending + +Magics distinct from BSHK/BSHR/BCGB/BCGR/BSCP/BSCR. + +### Validation (3090-ai.foxhop.net, RTX 3090) + +Byte-identity vs Python `sorted()` PASS at n ∈ {32, 1k, 100k, 1M, 10M}. +Bench table from `--bench` mode: + +| n | kernel ms | Gkeys/s | +|-----------|-----------|---------| +| 100,000 | 0.142 | 0.706 | +| 1,000,000 | 0.265 | 3.767 | +| 10,000,000| 1.817 | 5.504 | + +Reference catalog number: 1.4 Gkeys/s on Titan. 3090 lands ~4x that at +saturation, which matches CUB's published Ampere scaling. Kernel time +at small N reflects fixed launch overhead, not algorithm cost — Gkeys/s +only stabilizes past n ≈ 1M. + +### 4090 host (ai.foxhop.net) — capacity constraint, not a form-G defect + +`./radix-sort --binary` OOMs on cudaMalloc for any N when all four +gpu-worker daemons are co-resident: the secp256k1 daemon allocates +its full scratch footprint at startup, leaving 47 MiB free of the +4090's 24 GiB. This is a pre-existing constraint of the ai box's +daemon mix; radix-sort itself built clean & registered with the worker +log line `gpu-worker: ready cuda-radix-sort <- ./radix-sort`. + +Standalone validation on the 4090 will run once the secp daemon +shrinks its preallocation OR we add a `RADIX_SORT_ONLY=1` worker mode. +Out of scope for form-G day-1; tracked here as a follow-up. + +### Files + +- `radix-sort.cu` — daemon + binary mode, mirroring shake256-fanout.cu +- `Makefile` — `radix-sort`, `radix-sort-test`, `radix-sort-bench` targets +- `test_radix_sort_known_answers.py` — sweep + bench harness +- `gpu-worker.lsp` — `handle-binary-sort` + `BSRT` magic dispatch + + `maybe-register-daemon!` of `cuda-radix-sort` diff --git a/examples/cuda-fanout/radix-sort.cu b/examples/cuda-fanout/radix-sort.cu new file mode 100644 index 0000000..5c2ec2b --- /dev/null +++ b/examples/cuda-fanout/radix-sort.cu @@ -0,0 +1,252 @@ +/* radix-sort.cu — bend form G: batched 64-bit key sort via NVIDIA CUB. + * + * Wraps cub::DeviceRadixSort::SortKeys on a u64 key stream. One-shot + * sort kernel; the dispatcher reserves op_id slots for descending + * order, u32 keys, & key-value pairs but only 0x01 (sort-u64-asc) is + * wired day-1. + * + * CATALOG.md form G entry: 1.4 G keys/sec on Titan; 20-50x over CPU + * merge sort; 257x over Intel Xeon Phi for scan. + * + * Wire (BSRT input): + * "BSRT" (4 B magic) + * u32 op_id (0x01 = sort-u64-asc) + * u32 n_items + * u64[n_items] (little-endian keys) + * + * Wire (BSRR output): + * "BSRR" (4 B magic) + * u32 status (0 = ok) + * u32 n_items + * u64[n_items] (sorted ascending) + * + * Reserved op slots (NOT day-1): + * 0x02 sort-u64-desc + * 0x03 sort-u32-asc / desc + * 0x04 sort-key-value pairs + * + * Build: + * make radix-sort + * + * License: AGPLv3. CUB is BSD-3-Clause (NVIDIA), header-only, so the + * binary is a thin wrapper around their device sort. + */ + +#include +#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)); \ + return -1; \ + } \ +} while (0) + +/* Op codes. Only OP_SORT_U64_ASC is wired today; the rest hold their + * slot so a future build can grow the dispatch without renumbering. */ +enum { + OP_SORT_U64_ASC = 0x01, + OP_SORT_U64_DESC = 0x02, + OP_SORT_U32 = 0x03, /* asc/desc selected by low bit later */ + OP_SORT_KV = 0x04, /* key-value pairs */ +}; + +/* ── small endian helpers ─────────────────────────────────────────── */ + +static uint32_t rd32_le(const uint8_t *p) { + return (uint32_t)p[0] | ((uint32_t)p[1] << 8) + | ((uint32_t)p[2] << 16) | ((uint32_t)p[3] << 24); +} + +static void wr32_le(uint8_t *p, uint32_t v) { + p[0] = v & 0xff; p[1] = (v >> 8) & 0xff; + p[2] = (v >> 16) & 0xff; p[3] = (v >> 24) & 0xff; +} + +static int read_file_to_buffer(const char *path, uint8_t **out_buf, size_t *out_len) { + FILE *f = fopen(path, "rb"); + if (!f) { fprintf(stderr, "open %s failed\n", 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 (!buf) { fclose(f); return -1; } + if (fread(buf, 1, sz, f) != (size_t)sz) { + free(buf); fclose(f); return -1; + } + fclose(f); + *out_buf = buf; + *out_len = sz; + return 0; +} + +static int write_buffer_to_file(const char *path, const uint8_t *buf, size_t len) { + FILE *f = fopen(path, "wb"); + if (!f) { fprintf(stderr, "create %s failed\n", path); return -1; } + if (fwrite(buf, 1, len, f) != len) { fclose(f); return -1; } + fclose(f); + return 0; +} + +/* ── core: BSRT in → BSRR out ─────────────────────────────────────── */ + +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); + + uint8_t *in_buf = NULL; + size_t in_len = 0; + if (read_file_to_buffer(in_path, &in_buf, &in_len) != 0) return -1; + + if (in_len < 12 || memcmp(in_buf, "BSRT", 4) != 0) { + fprintf(stderr, "bad BSRT magic / short input (%zu bytes)\n", in_len); + free(in_buf); return -1; + } + uint32_t op_id = rd32_le(in_buf + 4); + uint32_t n = rd32_le(in_buf + 8); + + if (op_id != OP_SORT_U64_ASC) { + fprintf(stderr, "unsupported op_id 0x%02x (only 0x01 sort-u64-asc wired)\n", + op_id); + free(in_buf); return -1; + } + + size_t expected = 12 + (size_t)n * sizeof(uint64_t); + if (in_len != expected) { + fprintf(stderr, "BSRT payload size %zu != expected %zu (op=0x%02x n=%u)\n", + in_len, expected, op_id, n); + free(in_buf); return -1; + } + + const uint64_t *h_keys_in = (const uint64_t*)(in_buf + 12); + + /* Allocate device buffers (input + output ping-pong CUB needs). */ + uint64_t *d_keys_in = NULL; + uint64_t *d_keys_out = NULL; + size_t n_bytes = (size_t)n * sizeof(uint64_t); + /* Always allocate at least 1 byte to keep cudaMalloc happy at n=0. */ + size_t alloc_bytes = n_bytes ? n_bytes : 1; + CUDA_CHECK(cudaMalloc(&d_keys_in, alloc_bytes)); + CUDA_CHECK(cudaMalloc(&d_keys_out, alloc_bytes)); + if (n > 0) { + CUDA_CHECK(cudaMemcpy(d_keys_in, h_keys_in, n_bytes, + cudaMemcpyHostToDevice)); + } + clock_gettime(CLOCK_MONOTONIC, &t1); + + /* CUB DeviceRadixSort: two-call pattern. First call with + * d_temp_storage=NULL returns the temp-storage size; second call + * does the actual sort. End-bit defaults to sizeof(KeyT)*8 (64). */ + void *d_temp = NULL; + size_t temp_bytes = 0; + cub::DeviceRadixSort::SortKeys( + d_temp, temp_bytes, + d_keys_in, d_keys_out, + (int)n); + CUDA_CHECK(cudaGetLastError()); + if (temp_bytes == 0) temp_bytes = 1; + CUDA_CHECK(cudaMalloc(&d_temp, temp_bytes)); + clock_gettime(CLOCK_MONOTONIC, &t2); + + cub::DeviceRadixSort::SortKeys( + d_temp, temp_bytes, + d_keys_in, d_keys_out, + (int)n); + CUDA_CHECK(cudaGetLastError()); + CUDA_CHECK(cudaDeviceSynchronize()); + clock_gettime(CLOCK_MONOTONIC, &t3); + + /* Pull sorted keys back, frame BSRR. */ + size_t out_size = 12 + n_bytes; + uint8_t *out_buf = (uint8_t*)malloc(out_size ? out_size : 1); + if (!out_buf) { + cudaFree(d_keys_in); cudaFree(d_keys_out); cudaFree(d_temp); + free(in_buf); return -1; + } + memcpy(out_buf, "BSRR", 4); + wr32_le(out_buf + 4, 0); /* status = ok */ + wr32_le(out_buf + 8, n); + if (n > 0) { + CUDA_CHECK(cudaMemcpy(out_buf + 12, d_keys_out, n_bytes, + cudaMemcpyDeviceToHost)); + } + clock_gettime(CLOCK_MONOTONIC, &t4); + + cudaFree(d_keys_in); + cudaFree(d_keys_out); + cudaFree(d_temp); + free(in_buf); + + int wr = write_buffer_to_file(out_path, out_buf, out_size); + free(out_buf); + + double ld_ms = (t1.tv_sec - t0.tv_sec) * 1e3 + (t1.tv_nsec - t0.tv_nsec) / 1e6; + double probe_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; + double gkeys = (kern_ms > 0.0) ? ((double)n / kern_ms / 1e6) : 0.0; + fprintf(stderr, + "radix-sort op=0x%02x n=%u load=%.2f probe=%.2f kernel=%.3f write=%.2f ms (%.3f Gkeys/s)\n", + op_id, n, ld_ms, probe_ms, kern_ms, write_ms, gkeys); + return wr; +} + +/* ── main: --binary one-shot, --daemon long-lived ─────────────────── */ + +int main(int argc, char **argv) { + if (argc >= 2 && !strcmp(argv[1], "--daemon")) { + int dev_count = 0; + cudaGetDeviceCount(&dev_count); + if (dev_count == 0) { + fprintf(stderr, "no CUDA devices\n"); + return 1; + } + cudaFree(0); /* eager 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-bin ", 12)) { + char *args = line + 12; + char *sep = strchr(args, ' '); + if (!sep) { + fprintf(stdout, "error bad-args\n"); fflush(stdout); + continue; + } + *sep = 0; + char *in_path = args; + char *out_path = sep + 1; + int r = process_one_bin(in_path, out_path); + fprintf(stdout, r == 0 ? "done %s\n" : "error process-failed\n", + out_path); + fflush(stdout); + } else if (line[0]) { + fprintf(stdout, "error unknown-command\n"); fflush(stdout); + } + } + return 0; + } + + if (argc >= 4 && !strcmp(argv[1], "--binary")) { + return process_one_bin(argv[2], argv[3]); + } + + fprintf(stderr, + "usage: %s --binary \n" + " %s --daemon (read commands on stdin)\n", + argv[0], argv[0]); + return 1; +} diff --git a/examples/cuda-fanout/test_radix_sort_known_answers.py b/examples/cuda-fanout/test_radix_sort_known_answers.py new file mode 100644 index 0000000..ab217ea --- /dev/null +++ b/examples/cuda-fanout/test_radix_sort_known_answers.py @@ -0,0 +1,184 @@ +"""test_radix_sort_known_answers.py — validation for radix-sort. + +Pipes BSRT-format u64 key arrays at the CUB radix-sort worker via +--binary mode, parses the BSRR response, & compares the byte stream +against Python's sorted() reference. + +Coverage: op_id 0x01 (sort-u64-asc). Sweeps N across small, medium, +& large workloads so the test catches both edge cases (tiny N) and +sustained-throughput regressions (10 M keys). + +Usage: + python3 -u test_radix_sort_known_answers.py [worker] # default ./radix-sort + python3 -u test_radix_sort_known_answers.py ./radix-sort --bench + python3 -u test_radix_sort_known_answers.py ./radix-sort --n 1000000 +""" +import argparse +import os +import random +import re +import subprocess +import sys +import tempfile +import time + +OP_SORT_U64_ASC = 0x01 + + +def build_bsrt(op_id, keys): + """Frame an u64 key list into a BSRT payload.""" + n = len(keys) + parts = [ + b"BSRT", + op_id.to_bytes(4, "little"), + n.to_bytes(4, "little"), + ] + for k in keys: + parts.append(int(k).to_bytes(8, "little")) + return b"".join(parts) + + +def parse_bsrr(blob, n_expected): + if len(blob) < 12 or blob[:4] != b"BSRR": + raise RuntimeError(f"bad BSRR magic / short: {blob[:40]!r}") + status = int.from_bytes(blob[4:8], "little") + n = int.from_bytes(blob[8:12], "little") + if status != 0: + raise RuntimeError(f"worker returned status {status}") + if n != n_expected: + raise RuntimeError(f"BSRR n {n} != expected {n_expected}") + expected_len = 12 + n * 8 + if len(blob) != expected_len: + raise RuntimeError( + f"BSRR size {len(blob)} != expected {expected_len} (n={n})" + ) + body = blob[12:] + return [int.from_bytes(body[i*8:(i+1)*8], "little") for i in range(n)] + + +KERNEL_RE = re.compile( + r"kernel=([\d.]+)\s+write=[\d.]+\s+ms\s+\(([\d.]+)\s+Gkeys/s\)" +) + + +def run_worker_binary(worker, in_blob): + with tempfile.NamedTemporaryFile(delete=False, suffix=".bsrt") as fin: + fin.write(in_blob) + in_path = fin.name + out_path = in_path.replace(".bsrt", ".bsrr") + try: + result = subprocess.run( + [worker, "--binary", in_path, out_path], + capture_output=True, timeout=300, + ) + if result.returncode != 0: + raise RuntimeError( + f"worker exited {result.returncode}: {result.stderr.decode()[:400]}" + ) + with open(out_path, "rb") as fout: + return fout.read(), result.stderr.decode() + finally: + for p in (in_path, out_path): + try: + os.unlink(p) + except OSError: + pass + + +def parse_kernel_telemetry(stderr_text): + """Extract (kernel_ms, gkeys_per_sec) from the worker's stderr line.""" + m = KERNEL_RE.search(stderr_text) + if not m: + return None, None + return float(m.group(1)), float(m.group(2)) + + +def test_n(worker, n, seed=0xC0FFEE57C0DEC0DE): + """Sort N random u64s, compare byte-for-byte against Python sorted().""" + rng = random.Random(seed ^ n) + keys = [rng.randrange(0, 1 << 64) for _ in range(n)] + expected = sorted(keys) + in_blob = build_bsrt(OP_SORT_U64_ASC, keys) + + t_host = time.perf_counter() + out_blob, stderr = run_worker_binary(worker, in_blob) + wall_ms = (time.perf_counter() - t_host) * 1000.0 + got = parse_bsrr(out_blob, n) + + kernel_ms, gkeys = parse_kernel_telemetry(stderr) + # Byte-identity check: rebuild canonical bytes from sorted() and + # compare to the worker's body (skip the 12-byte BSRR header). + expected_bytes = b"".join(int(k).to_bytes(8, "little") for k in expected) + actual_bytes = out_blob[12:] + ok = (expected_bytes == actual_bytes) + + if not ok: + # Locate first mismatch for the report. + first = None + for i in range(n): + if got[i] != expected[i]: + first = (i, expected[i], got[i]) + break + kernel_label = f"{kernel_ms:.3f}" if kernel_ms is not None else "n/a" + print(f" FAIL sort-u64-asc n={n:<10} wall={wall_ms:.1f}ms kernel={kernel_label}ms") + if first is not None: + i, exp, gv = first + print(f" first mismatch @ i={i}: expected=0x{exp:016x} got=0x{gv:016x}") + return False, kernel_ms, gkeys + + extra = "" + if kernel_ms is not None and gkeys is not None: + extra = f" kernel={kernel_ms:.3f}ms {gkeys:.3f} Gkeys/s" + print(f" PASS sort-u64-asc n={n:<10} wall={wall_ms:.1f}ms{extra}") + return True, kernel_ms, gkeys + + +DEFAULT_NS = [32, 1_000, 100_000, 1_000_000, 10_000_000] +BENCH_NS = [100_000, 1_000_000, 10_000_000] + + +def main(): + ap = argparse.ArgumentParser() + ap.add_argument("worker", nargs="?", default="./radix-sort") + ap.add_argument("--n", type=int, default=None, + help="single N override (default: sweep)") + ap.add_argument("--bench", action="store_true", + help="report kernel ms / Gkeys/s table at n=100k/1M/10M") + args = ap.parse_args() + + if not os.path.exists(args.worker): + print(f"worker not found: {args.worker}") + sys.exit(1) + + if args.n: + ns = [args.n] + elif args.bench: + ns = BENCH_NS + else: + ns = DEFAULT_NS + + all_pass = True + bench_rows = [] + for n in ns: + ok, kernel_ms, gkeys = test_n(args.worker, n) + all_pass = all_pass and ok + bench_rows.append((n, kernel_ms, gkeys)) + + if args.bench: + print("\nbench summary") + print(f" {'n':>10} {'kernel_ms':>12} {'Gkeys/s':>10}") + for n, kernel_ms, gkeys in bench_rows: + km = f"{kernel_ms:.3f}" if kernel_ms is not None else "n/a" + gk = f"{gkeys:.3f}" if gkeys is not None else "n/a" + print(f" {n:>10} {km:>12} {gk:>10}") + + if all_pass: + print("\nALL PASS") + sys.exit(0) + else: + print("\nFAILED") + sys.exit(1) + + +if __name__ == "__main__": + main() diff --git a/www/bend.html b/www/bend.html index 17e158e..a2c3c5b 100644 --- a/www/bend.html +++ b/www/bend.html @@ -125,6 +125,12 @@ make gpu-worker LUMBDA=asm # smallest footprint 13.83 Mkeys/s @ n=1M (~309× coincurve CPU; windowed-G ladder w=4) BSCP binary: scalars + base-point → BSCR points + + cuda-radix-sort + RTX 3090 + 5.50 Gkeys/s @ n=10M (kernel 1.82 ms); 3.77 Gkeys/s @ n=1M; CUB DeviceRadixSort u64 ascending + BSRT binary: op_id + n + u64[n] → BSRR sorted u64[n] + @@ -138,7 +144,7 @@ make gpu-worker LUMBDA=asm # smallest footprint cuda-clifford-stabilizer186× over Stim (CPU SOTA) on equivalence-checkingSTABSimSTABSim · Qimax cuda-bernstein-yang-inv3–10× per inversion over Fermat on CPU; novel territory on GPU—safegcd · Jumping cuda-ntt-polyup to 123× over CPU; 21× on RTX 3070RTX 3070NTTSuite · FHE NTT - cuda-radix-sort1.4 G keys/sec; 20–50× over CPU merge sort; 257× vs Xeon Phi for scanTitanCUB · Onesweep + cuda-radix-sort1.4 G keys/sec; 20–50× over CPU merge sort; 257× vs Xeon Phi for scan(promoted — see Live)CUB · Onesweep