bend form G — cuda-radix-sort lands live (CUB DeviceRadixSort u64)
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.
This commit is contained in:
parent
948bdf0e41
commit
05fa922d42
7 changed files with 553 additions and 2 deletions
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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.
|
||||
|
|
|
|||
|
|
@ -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)
|
||||
|
|
|
|||
52
examples/cuda-fanout/plans/form-G-progress.md
Normal file
52
examples/cuda-fanout/plans/form-G-progress.md
Normal file
|
|
@ -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`
|
||||
252
examples/cuda-fanout/radix-sort.cu
Normal file
252
examples/cuda-fanout/radix-sort.cu
Normal file
|
|
@ -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 <cuda_runtime.h>
|
||||
#include <cub/device/device_radix_sort.cuh>
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <string.h>
|
||||
#include <stdint.h>
|
||||
#include <time.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)); \
|
||||
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 <input.bin> <output.bin>\n"
|
||||
" %s --daemon (read commands on stdin)\n",
|
||||
argv[0], argv[0]);
|
||||
return 1;
|
||||
}
|
||||
184
examples/cuda-fanout/test_radix_sort_known_answers.py
Normal file
184
examples/cuda-fanout/test_radix_sort_known_answers.py
Normal file
|
|
@ -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()
|
||||
|
|
@ -125,6 +125,12 @@ make gpu-worker LUMBDA=asm # smallest footprint</code></pre>
|
|||
<td>13.83 Mkeys/s @ n=1M (~309× coincurve CPU; windowed-G ladder w=4)</td>
|
||||
<td><code>BSCP</code> binary: scalars + base-point → <code>BSCR</code> points</td>
|
||||
</tr>
|
||||
<tr>
|
||||
<td><code>cuda-radix-sort</code></td>
|
||||
<td>RTX 3090</td>
|
||||
<td>5.50 Gkeys/s @ n=10M (kernel 1.82 ms); 3.77 Gkeys/s @ n=1M; CUB DeviceRadixSort u64 ascending</td>
|
||||
<td><code>BSRT</code> binary: op_id + n + u64[n] → <code>BSRR</code> sorted u64[n]</td>
|
||||
</tr>
|
||||
</tbody>
|
||||
</table>
|
||||
|
||||
|
|
@ -138,7 +144,7 @@ make gpu-worker LUMBDA=asm # smallest footprint</code></pre>
|
|||
<tr><td><code>cuda-clifford-stabilizer</code></td><td>186× over Stim (CPU SOTA) on equivalence-checking</td><td>STABSim</td><td><a href="https://arxiv.org/abs/2507.03092">STABSim</a> · <a href="https://arxiv.org/abs/2505.03307">Qimax</a></td></tr>
|
||||
<tr><td><code>cuda-bernstein-yang-inv</code></td><td>3–10× per inversion over Fermat on CPU; novel territory on GPU</td><td>—</td><td><a href="https://eprint.iacr.org/2019/266">safegcd</a> · <a href="https://eprint.iacr.org/2024/644">Jumping</a></td></tr>
|
||||
<tr><td><code>cuda-ntt-poly</code></td><td>up to 123× over CPU; 21× on RTX 3070</td><td>RTX 3070</td><td><a href="https://arxiv.org/pdf/2405.11353">NTTSuite</a> · <a href="https://eprint.iacr.org/2021/124.pdf">FHE NTT</a></td></tr>
|
||||
<tr><td><code>cuda-radix-sort</code></td><td>1.4 G keys/sec; 20–50× over CPU merge sort; 257× vs Xeon Phi for scan</td><td>Titan</td><td><a href="https://github.com/NVIDIA/cub">CUB</a> · <a href="https://gpuopen.com/learn/boosting_gpu_radix_sort/">Onesweep</a></td></tr>
|
||||
<tr><td><code>cuda-radix-sort</code></td><td>1.4 G keys/sec; 20–50× over CPU merge sort; 257× vs Xeon Phi for scan</td><td>(promoted — see Live)</td><td><a href="https://github.com/NVIDIA/cub">CUB</a> · <a href="https://gpuopen.com/learn/boosting_gpu_radix_sort/">Onesweep</a></td></tr>
|
||||
</tbody>
|
||||
</table>
|
||||
|
||||
|
|
|
|||
Loading…
Add table
Add a link
Reference in a new issue