diff --git a/examples/cuda-fanout/CATALOG.md b/examples/cuda-fanout/CATALOG.md index 57c33c1..5572d81 100644 --- a/examples/cuda-fanout/CATALOG.md +++ b/examples/cuda-fanout/CATALOG.md @@ -25,6 +25,7 @@ form-status column says `planned` until numbers exist. | `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-blake3-tree` | `blake3-fanout` | 3090 | **32.5 GB/s @ 1M × 64 B** (kernel 1.97 ms); 20.4 GB/s @ 1k × 1 MB (kernel 51.5 ms); 2.68 GB/s @ 100k × 64 B; byte-identical vs BLAKE3 reference spec across single-chunk + multi-chunk paths at n∈{32, 1k, 10k, 100k, 1M}; one CUDA thread per input walks Merkle chunk tree on-device | `BSB3` binary: out_bytes + n + (u32 len + bytes) per input → `BSR3` n + out_bytes + digests | `cuda-sim-ops-bin` numbers are humbling on purpose: kickmix is conditional-op heavy & branch-divergent. This catalog exists so @@ -125,7 +126,7 @@ run on our fleet yet. | cuda-dilithium-pqsig | 57.7x keygen+sign+verify | RTX 3090 Ti | HIGH (PQ TLS migration) | | cuda-mc-options-pricing | 25-152x (barrier-call kernel) | Tesla C1060 / modern | LOW (calibration form) | | cuda-cuFFT-batched-1D | 8-32x vs MKL; tcFFT 1.1-3.2x over cuFFT| V100 / A100 | MEDIUM (punters-cc spectrograms) | -| cuda-blake3-tree | ~5-20x (tree mode) | Blaze-3 CUDA | HIGH (content-addressed portals) | +| cuda-blake3-tree | **LIVE** — 32.5 GB/s @ 1M × 64 B; 20.4 GB/s @ 1k × 1 MB on 3090; see Live forms table | RTX 3090 | HIGH (content-addressed portals) | | cuda-bloom-filter-modern | ~6x CPU; 3.4B inserts/s | B200 / Perlmutter | HIGH (foxhop pruning, undefect known-set) | | cuda-gemm-batched-FP8 | 4.8x FP8 over A100; 716 TFLOPS H100 | H100 SXM | LOW (calibration / lattice PQC) | | cuda-batched-matrix-inverse | 4.3-16.8x vs MAGMA | P100 | LOW (LA verifiers on circuits) | diff --git a/examples/cuda-fanout/Makefile b/examples/cuda-fanout/Makefile index 27e3644..3a52d7d 100644 --- a/examples/cuda-fanout/Makefile +++ b/examples/cuda-fanout/Makefile @@ -16,11 +16,28 @@ 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 radix-sort +all: shake256-fanout cgbn-batch-worker secp256k1-batch-mul radix-sort blake3-fanout shake256-fanout: shake256-fanout.cu $(NVCC) $(NVCCFLAGS) -o $@ $< +# Wave-2 form `cuda-blake3-tree` — BLAKE3 batched-hash fan-out. +# One CUDA thread per input, internal Merkle tree walked on-device. +# Compression primitives vendored from Blaze-3/BLAKE3-gpu (MIT, see +# vendor/blake3-gpu/{LICENSE,NOTICE}); wire wrapper AGPLv3. +# Wire: BSB3 request → BSR3 response, distinct from BSHK/BCGB/BSCP/BSRT. +blake3-fanout: blake3-fanout.cu vendor/blake3-gpu/blake3_device.cuh + $(NVCC) $(NVCCFLAGS) -o $@ $< + +blake3-test: blake3-fanout test_blake3_known_answers.py + python3 -u test_blake3_known_answers.py ./blake3-fanout --quick + +blake3-test-full: blake3-fanout test_blake3_known_answers.py + python3 -u test_blake3_known_answers.py ./blake3-fanout + +blake3-bench: blake3-fanout test_blake3_known_answers.py + python3 -u test_blake3_known_answers.py ./blake3-fanout --n 1000000 --mode fixed-64 + # 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. diff --git a/examples/cuda-fanout/blake3-fanout.cu b/examples/cuda-fanout/blake3-fanout.cu new file mode 100644 index 0000000..ab1b1e7 --- /dev/null +++ b/examples/cuda-fanout/blake3-fanout.cu @@ -0,0 +1,545 @@ +/* blake3-fanout.cu — BLAKE3 batched-hash CUDA worker, Wave-2 form + * `cuda-blake3-tree`. Sibling of shake256-fanout.cu — same wire + * shape, same daemon contract, distinct binary magic (BSB3 / BSR3). + * + * One CUDA thread per input. Each thread walks the BLAKE3 chunk-tree + * serially: per-chunk compression, then parent merges, then a root + * compression. Captures both single-chunk (input <= 1024 B) and + * multi-chunk paths. + * + * Wire format mirrors shake256-fanout exactly. Binary mode: + * request: "BSB3" | u32 out_bytes | u32 n_items | for each: + * u32 len | len bytes + * response: "BSR3" | u32 n_items | u32 out_bytes | + * for each: out_bytes + * + * Daemon protocol on stdin (one line per command): + * process — S-expr in, S-expr out + * process-bin — binary in, binary out + * quit + * + * Stdout response: "done " / "error " / "bye". + * + * License: AGPLv3. BLAKE3 __device__ primitives vendored from + * Blaze-3/BLAKE3-gpu (MIT, see vendor/blake3-gpu/{LICENSE,NOTICE}). + * Constants and round function track the BLAKE3 reference spec + * (https://github.com/BLAKE3-team/BLAKE3 — CC0 / Apache-2.0). + */ + +#include +#include +#include +#include +#include +#include + +#include "vendor/blake3-gpu/blake3_device.cuh" + +using blake3_gpu::u8; +using blake3_gpu::u32; +using blake3_gpu::u64; +using blake3_gpu::IV; +using blake3_gpu::CHUNK_START; +using blake3_gpu::CHUNK_END; +using blake3_gpu::PARENT; +using blake3_gpu::ROOT; +using blake3_gpu::BLOCK_LEN; +using blake3_gpu::CHUNK_LEN; +using blake3_gpu::OUT_LEN; +using blake3_gpu::compress; +using blake3_gpu::words_from_bytes_le; + +#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) + +/* ─── per-thread BLAKE3 over one input ──────────────────────────────── */ + +/* Produce the 8-word chaining value of a *complete* (1024 B) chunk at + * `counter`. Used only for chunks 0..N-2 of multi-chunk inputs. */ +__device__ static void chunk_cv_complete(const u8 *input, u64 counter, + u32 *out_cv8) +{ + u32 cv[8]; + for (int i = 0; i < 8; i++) cv[i] = IV[i]; + + u32 block_words[16]; + u32 state[16]; + + /* CHUNK_LEN = 16 * BLOCK_LEN — exactly 16 compressions per full chunk. */ + for (u32 b = 0; b < (CHUNK_LEN / BLOCK_LEN); b++) { + words_from_bytes_le(input + (size_t)b * BLOCK_LEN, BLOCK_LEN, block_words); + u32 flags = 0; + if (b == 0) flags |= CHUNK_START; + if (b == (CHUNK_LEN / BLOCK_LEN) - 1) flags |= CHUNK_END; + compress(cv, block_words, counter, BLOCK_LEN, flags, state); + for (int i = 0; i < 8; i++) cv[i] = state[i]; + } + for (int i = 0; i < 8; i++) out_cv8[i] = cv[i]; +} + +/* Build the "pending Output" for the final chunk (or the only chunk). The + * output is captured as (input_cv, block_words, counter, block_len, flags) + * — we do not run the final compress yet; finalize-time decides ROOT. + * + * `input` / `len` describe the chunk's bytes; `counter` is chunk index; + * `is_root_chunk` is true ONLY for single-chunk inputs (no parents will + * be applied later — we still defer the final compress so that the + * caller adds ROOT during finalize). */ +struct Pending { + u32 input_cv[8]; + u32 block_words[16]; + u64 counter; + u32 block_len; + u32 flags; +}; + +__device__ static void chunk_pending(const u8 *input, u32 len, u64 counter, + Pending &p) +{ + u32 cv[8]; + for (int i = 0; i < 8; i++) cv[i] = IV[i]; + + u32 block_words[16]; + u32 state[16]; + + u32 offset = 0; + u32 blocks_compressed = 0; + + /* Compress every full preceding block. Handle empty input (len==0): we + * enter the post-loop with blocks_compressed==0 and emit one final + * 0-byte block with CHUNK_START | CHUNK_END. */ + while (offset + BLOCK_LEN < len) { + words_from_bytes_le(input + offset, BLOCK_LEN, block_words); + u32 flags = 0; + if (blocks_compressed == 0) flags |= CHUNK_START; + compress(cv, block_words, counter, BLOCK_LEN, flags, state); + for (int i = 0; i < 8; i++) cv[i] = state[i]; + blocks_compressed++; + offset += BLOCK_LEN; + } + + /* Final block: bytes [offset, len). Pad inside words_from_bytes_le. */ + u32 final_len = len - offset; + words_from_bytes_le(input + offset, final_len, p.block_words); + for (int i = 0; i < 8; i++) p.input_cv[i] = cv[i]; + p.counter = counter; + p.block_len = final_len; + p.flags = CHUNK_END; + if (blocks_compressed == 0) p.flags |= CHUNK_START; +} + +/* BLAKE3 tree walk for a single input. Stack depth ≤ 54 per spec + * (2^54 * 1024 B = 2^64). Cap at 56 for slack. Each entry = 8 u32 = 32 B. */ +__device__ static void blake3_hash_one(const u8 *input, u32 len, u8 *out32) +{ + constexpr int STACK_MAX = 56; + u32 stack_cv[STACK_MAX][8]; + int stack_depth = 0; + + /* Walk all chunks except the last. Each complete chunk produces a CV + * pushed onto the stack with the trailing-zeros merge rule. */ + u64 chunk_idx = 0; + u32 offset = 0; + /* `last_chunk_start` = offset of the last (possibly partial) chunk. */ + u32 last_chunk_start; + if (len == 0) { + last_chunk_start = 0; + } else if (len % CHUNK_LEN == 0) { + last_chunk_start = len - CHUNK_LEN; + } else { + last_chunk_start = (len / CHUNK_LEN) * CHUNK_LEN; + } + + while (offset < last_chunk_start) { + u32 cv8[8]; + chunk_cv_complete(input + offset, chunk_idx, cv8); + + /* add_chunk_chaining_value: merge while total_chunks (chunk_idx+1) + * has trailing zero bits. Each merge pops one CV, parents with the + * incoming cv8, and the result becomes the new cv8. */ + u64 total_chunks = chunk_idx + 1; + while ((total_chunks & 1ULL) == 0ULL) { + u32 left[8]; + for (int i = 0; i < 8; i++) left[i] = stack_cv[stack_depth - 1][i]; + stack_depth--; + + u32 block_words[16]; + for (int i = 0; i < 8; i++) block_words[i] = left[i]; + for (int i = 0; i < 8; i++) block_words[8 + i] = cv8[i]; + u32 parent_cv[8]; + for (int i = 0; i < 8; i++) parent_cv[i] = IV[i]; + u32 state[16]; + compress(parent_cv, block_words, 0ULL, BLOCK_LEN, PARENT, state); + for (int i = 0; i < 8; i++) cv8[i] = state[i]; + total_chunks >>= 1; + } + /* push cv8 */ + for (int i = 0; i < 8; i++) stack_cv[stack_depth][i] = cv8[i]; + stack_depth++; + + chunk_idx++; + offset += CHUNK_LEN; + } + + /* Build the pending Output for the final chunk. */ + Pending p; + chunk_pending(input + last_chunk_start, len - last_chunk_start, chunk_idx, p); + + /* Finalize: fold the stack into Output right-to-left. Each fold computes + * the chaining value of the current Output (non-root) and uses it as + * the right child of a parent Output. The final Output (no more stack + * entries) gets ROOT applied when its compress runs. */ + while (stack_depth > 0) { + /* compute current Output's chaining value (non-root) */ + u32 child_cv[8]; + { + u32 state[16]; + compress(p.input_cv, p.block_words, p.counter, p.block_len, p.flags, state); + for (int i = 0; i < 8; i++) child_cv[i] = state[i]; + } + /* pop left from stack, build parent Output */ + u32 left[8]; + for (int i = 0; i < 8; i++) left[i] = stack_cv[stack_depth - 1][i]; + stack_depth--; + for (int i = 0; i < 8; i++) p.block_words[i] = left[i]; + for (int i = 0; i < 8; i++) p.block_words[8 + i] = child_cv[i]; + for (int i = 0; i < 8; i++) p.input_cv[i] = IV[i]; + p.counter = 0ULL; + p.block_len = BLOCK_LEN; + p.flags = PARENT; + } + + /* Apply ROOT and compress. First 32 B of root output = state[0..8] LE. */ + u32 state[16]; + compress(p.input_cv, p.block_words, 0ULL, p.block_len, p.flags | ROOT, state); + for (int i = 0; i < 8; i++) { + u32 w = state[i]; + out32[4 * i + 0] = (u8)(w & 0xFF); + out32[4 * i + 1] = (u8)((w >> 8) & 0xFF); + out32[4 * i + 2] = (u8)((w >> 16) & 0xFF); + out32[4 * i + 3] = (u8)((w >> 24) & 0xFF); + } +} + +__global__ void blake3_fanout_kernel(const u8 *inputs, + const u32 *offsets, + const u32 *lengths, + u8 *outputs, + u32 out_bytes, + u32 n) +{ + u32 t = blockIdx.x * blockDim.x + threadIdx.x; + if (t >= n) return; + /* We only realize the first 32 bytes of the root output. */ + u8 digest[32]; + blake3_hash_one(inputs + offsets[t], lengths[t], digest); + u32 to_copy = out_bytes < 32 ? out_bytes : 32; + for (u32 i = 0; i < to_copy; i++) { + outputs[(size_t)t * out_bytes + i] = digest[i]; + } + for (u32 i = 32; i < out_bytes; i++) { + /* out_bytes > 32: zero-pad. Future extension can XOF here. */ + outputs[(size_t)t * out_bytes + i] = 0; + } +} + +/* ─── 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; +} + +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; +} + +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) { + 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); /* forces 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; + } + 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); + continue; + } + *sep = 0; + char *out_path = sep + 1; + 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"); + 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]); + } + if (argc < 3) { + fprintf(stderr, + "usage: %s \n" + " %s --binary \n" + " %s --daemon (read commands on stdin)\n", + argv[0], argv[0], argv[0]); + return 1; + } + return process_one(argv[1], argv[2]); +} + +/* Binary format (mirrors shake256-fanout layout, distinct magic): + * in: "BSB3" | u32 out_bytes | u32 n | (u32 len | len bytes) × n + * out: "BSR3" | u32 n | u32 out_bytes | (out_bytes bytes) × n + */ +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); + + if (sz < 12 || memcmp(buf, "BSB3", 4) != 0) { + fprintf(stderr, "%s: bad magic (want BSB3)\n", in_path); return 1; + } + uint32_t out_bytes, n; + memcpy(&out_bytes, buf + 4, 4); + memcpy(&n, buf + 8, 4); + + uint32_t *offsets = (uint32_t*)malloc(n * sizeof(uint32_t)); + uint32_t *lengths = (uint32_t*)malloc(n * sizeof(uint32_t)); + size_t cur = 12; + size_t data_off = 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; + } + size_t total = data_off; + uint8_t *flat = (uint8_t*)malloc(total ? total : 1); + cur = 12; + 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; + blake3_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("BSR3", 1, 4, fo); + 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); + + size_t portal_len; + char *portal = slurp(in_path, &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", in_path); + return 1; + } + + 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); + + 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; + blake3_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 *f = fopen(out_path, "w"); + if (!f) { perror(out_path); return 1; } + fprintf(f, "(cuda-blake3-tree-result\n"); + fprintf(f, " (n %u)\n", n); + fprintf(f, " (output-bytes %u)\n", out_bytes); + fprintf(f, " (hashes\n"); + char *hex_buf = (char*)malloc(out_bytes * 2 + 1); + for (uint32_t i = 0; i < n; i++) { + bytes_to_hex(h_out + (size_t)i * out_bytes, out_bytes, hex_buf); + fprintf(f, " \"%s\"\n", hex_buf); + } + fprintf(f, " )\n"); + clock_gettime(CLOCK_MONOTONIC, &t4); + double ld_ms = (t1.tv_sec - t0.tv_sec) * 1e3 + (t1.tv_nsec - t0.tv_nsec) / 1e6; + double launch_ms = (t2.tv_sec - t1.tv_sec) * 1e3 + (t2.tv_nsec - t1.tv_nsec) / 1e6; + double kern_ms = (t3.tv_sec - t2.tv_sec) * 1e3 + (t3.tv_nsec - t2.tv_nsec) / 1e6; + double write_ms = (t4.tv_sec - t3.tv_sec) * 1e3 + (t4.tv_nsec - t3.tv_nsec) / 1e6; + fprintf(f, " (timing-ms (load %.3f) (launch %.3f) (kernel %.3f) (write %.3f)))\n", + ld_ms, launch_ms, kern_ms, write_ms); + fclose(f); + + cudaFree(d_in); cudaFree(d_off); cudaFree(d_len); cudaFree(d_out); + for (uint32_t i = 0; i < n; i++) free(hex_in[i]); + free(hex_in); free(offsets); free(lengths); free(flat); free(h_out); + free(hex_buf); free(portal); + return 0; +} diff --git a/examples/cuda-fanout/gpu-worker.lsp b/examples/cuda-fanout/gpu-worker.lsp index bebba04..8973d8a 100644 --- a/examples/cuda-fanout/gpu-worker.lsp +++ b/examples/cuda-fanout/gpu-worker.lsp @@ -56,6 +56,15 @@ (or (get-environment-variable "RADIX_SORT_WORKER") "./radix-sort")) +;; blake3-fanout — Wave-2 form `cuda-blake3-tree`. Per-input BLAKE3, +;; one CUDA thread walks the Merkle chunk tree for its input. Binary +;; wire uses BSB3 request / BSR3 response magic, distinct from all +;; other forms (BSHK/BCGB/BSCP/BSRT). Primitives vendored from +;; Blaze-3/BLAKE3-gpu (MIT); driver AGPLv3. +(define *binary-blake3-fanout* + (or (get-environment-variable "BLAKE3_FANOUT_WORKER") + "./blake3-fanout")) + (define (parse-port-arg args) (let loop ((rest args)) (cond @@ -364,6 +373,40 @@ (if (file-exists? out-path) (delete-file out-path)) (wire-send-raw client (string-append "BERR" (cdr status))))))))) +;; Binary wire for Wave-2 form `cuda-blake3-tree` (BLAKE3 batched fan-out). +;; Payload begins with "BSB3". The .cu daemon (blake3-fanout.cu) expects +;; the same magic + header in its on-disk binary file (matches the BCGB / +;; BSCP / BSRT pattern, not the BSHK strip-magic pattern), so we pass the +;; entire payload through verbatim. The daemon's response file already +;; begins with "BSR3"; pass through unchanged. +(define (handle-binary-blake3 client payload) + (let* ((daemon (cdr (assoc 'cuda-blake3-tree *daemons*))) + (in-path (gensym-path "/tmp/bend-blake3-in" ".bin")) + (out-path (gensym-path "/tmp/bend-blake3-out" ".bin")) + (t-start (current-time-ms))) + (write-binary-file in-path payload) + (display ";;; bend RECV cuda-blake3-tree 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-blake3-tree") + (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-blake3-tree 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))))))))) + ;; 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" @@ -422,6 +465,10 @@ (string=? (substring payload 0 4) "BSRT")) (handle-binary-sort client payload) (tcp-close client) #t) + ((and (>= (string-length payload) 4) + (string=? (substring payload 0 4) "BSB3")) + (handle-binary-blake3 client payload) + (tcp-close client) #t) (else (let* ((req (read-from-string payload)) (resp (handle-request req))) @@ -464,6 +511,7 @@ *binary-secp256k1-batch* *secp-daemon-extra-args*) (maybe-register-daemon! 'cuda-radix-sort *binary-radix-sort*) + (maybe-register-daemon! 'cuda-blake3-tree *binary-blake3-fanout*) (let ((server (tcp-listen port))) (cond ((eq? server #f) diff --git a/examples/cuda-fanout/test_blake3_known_answers.py b/examples/cuda-fanout/test_blake3_known_answers.py new file mode 100644 index 0000000..84dd408 --- /dev/null +++ b/examples/cuda-fanout/test_blake3_known_answers.py @@ -0,0 +1,368 @@ +"""test_blake3_known_answers.py — bend form `cuda-blake3-tree` validator. + +Pipes BSB3-format requests at our blake3-fanout worker via --binary mode, +parses the BSR3 response, and compares each digest against an in-process +BLAKE3 reference implementation translated directly from the upstream +reference spec (CC0 / Apache-2.0): + + https://github.com/BLAKE3-team/BLAKE3/blob/master/reference_impl/reference_impl.rs + +Reference covers chunk-tree walks for inputs of any length, including the +single-chunk fast path (<=1024 B) and multi-chunk Merkle paths. Both code +paths must match byte-for-byte. + +Wire (matches header doc in blake3-fanout.cu): + request: "BSB3" | u32 out_bytes | u32 n_items | (u32 len + bytes) * n + response: "BSR3" | u32 n_items | u32 out_bytes | (out_bytes bytes) * n + +Usage: + python3 -u test_blake3_known_answers.py [worker-path] + python3 -u test_blake3_known_answers.py ./blake3-fanout --n 1000 + python3 -u test_blake3_known_answers.py ./blake3-fanout --all # all sizes + +Exits 0 on PASS, 1 on FAIL. +""" +import argparse +import os +import random +import struct +import subprocess +import sys +import tempfile +import time + + +# ── BLAKE3 reference implementation (Python translation) ─────────────── +# Translated 1:1 from BLAKE3-team/BLAKE3 reference_impl/reference_impl.rs. + +OUT_LEN = 32 +KEY_LEN = 32 +BLOCK_LEN = 64 +CHUNK_LEN = 1024 + +CHUNK_START = 1 << 0 +CHUNK_END = 1 << 1 +PARENT = 1 << 2 +ROOT = 1 << 3 +KEYED_HASH = 1 << 4 + +IV = [ + 0x6A09E667, 0xBB67AE85, 0x3C6EF372, 0xA54FF53A, + 0x510E527F, 0x9B05688C, 0x1F83D9AB, 0x5BE0CD19, +] + +MSG_PERMUTATION = [2, 6, 3, 10, 7, 0, 4, 13, 1, 11, 12, 5, 9, 14, 15, 8] + +MASK32 = 0xFFFFFFFF + +def _rotr(x, n): + return ((x >> n) | (x << (32 - n))) & MASK32 + +def _g(state, a, b, c, d, mx, my): + state[a] = (state[a] + state[b] + mx) & MASK32 + state[d] = _rotr(state[d] ^ state[a], 16) + state[c] = (state[c] + state[d]) & MASK32 + state[b] = _rotr(state[b] ^ state[c], 12) + state[a] = (state[a] + state[b] + my) & MASK32 + state[d] = _rotr(state[d] ^ state[a], 8) + state[c] = (state[c] + state[d]) & MASK32 + state[b] = _rotr(state[b] ^ state[c], 7) + +def _round(state, m): + _g(state, 0, 4, 8, 12, m[0], m[1]) + _g(state, 1, 5, 9, 13, m[2], m[3]) + _g(state, 2, 6, 10, 14, m[4], m[5]) + _g(state, 3, 7, 11, 15, m[6], m[7]) + _g(state, 0, 5, 10, 15, m[8], m[9]) + _g(state, 1, 6, 11, 12, m[10], m[11]) + _g(state, 2, 7, 8, 13, m[12], m[13]) + _g(state, 3, 4, 9, 14, m[14], m[15]) + +def _permute(m): + p = [m[MSG_PERMUTATION[i]] for i in range(16)] + for i in range(16): + m[i] = p[i] + +def _compress(cv, block_words, counter, block_len, flags): + state = [ + cv[0], cv[1], cv[2], cv[3], + cv[4], cv[5], cv[6], cv[7], + IV[0], IV[1], IV[2], IV[3], + counter & MASK32, (counter >> 32) & MASK32, + block_len, flags, + ] + block = list(block_words) + for _ in range(6): + _round(state, block) + _permute(block) + _round(state, block) + for i in range(8): + state[i] ^= state[i + 8] + state[i+8] ^= cv[i] + return state + +def _words_from_le(byts): + n = len(byts) + pad = (-n) % 4 + if pad: + byts = bytes(byts) + b'\x00' * pad + return [int.from_bytes(byts[4*i:4*i+4], 'little') for i in range((n + pad) // 4)] + +def _le_bytes_from_words(words, n_bytes): + out = bytearray() + for w in words: + out.extend(int(w & MASK32).to_bytes(4, 'little')) + return bytes(out[:n_bytes]) + +class _Output: + def __init__(self, input_cv, block_words, counter, block_len, flags): + self.input_cv = list(input_cv) + self.block_words = list(block_words) + self.counter = counter + self.block_len = block_len + self.flags = flags + def chaining_value(self): + return _compress(self.input_cv, self.block_words, self.counter, + self.block_len, self.flags)[:8] + def root_output_bytes(self, n_bytes): + out = bytearray() + block_ctr = 0 + while len(out) < n_bytes: + words = _compress(self.input_cv, self.block_words, block_ctr, + self.block_len, self.flags | ROOT) + take = min(2 * OUT_LEN, n_bytes - len(out)) + out.extend(_le_bytes_from_words(words, take)) + block_ctr += 1 + return bytes(out[:n_bytes]) + +class _ChunkState: + def __init__(self, key, counter, flags): + self.cv = list(key) + self.counter = counter + self.buf = bytearray(BLOCK_LEN) + self.buf_len = 0 + self.blocks_compressed = 0 + self.flags = flags + def len(self): + return BLOCK_LEN * self.blocks_compressed + self.buf_len + def start_flag(self): + return CHUNK_START if self.blocks_compressed == 0 else 0 + def update(self, data): + i = 0 + while i < len(data): + if self.buf_len == BLOCK_LEN: + bw = _words_from_le(bytes(self.buf)) + self.cv = _compress(self.cv, bw, self.counter, BLOCK_LEN, + self.flags | self.start_flag())[:8] + self.blocks_compressed += 1 + self.buf = bytearray(BLOCK_LEN) + self.buf_len = 0 + take = min(BLOCK_LEN - self.buf_len, len(data) - i) + self.buf[self.buf_len:self.buf_len+take] = data[i:i+take] + self.buf_len += take + i += take + def output(self): + bw = _words_from_le(bytes(self.buf)) + return _Output(self.cv, bw, self.counter, self.buf_len, + self.flags | self.start_flag() | CHUNK_END) + +def _parent_output(left_cv, right_cv, key, flags): + bw = list(left_cv) + list(right_cv) + return _Output(list(key), bw, 0, BLOCK_LEN, PARENT | flags) + +def _parent_cv(left_cv, right_cv, key, flags): + return _parent_output(left_cv, right_cv, key, flags).chaining_value() + +class _Hasher: + def __init__(self, key=None, flags=0): + if key is None: key = IV + self.chunk_state = _ChunkState(key, 0, flags) + self.key = list(key) + self.cv_stack = [] + self.flags = flags + def _add_chunk_cv(self, new_cv, total_chunks): + while (total_chunks & 1) == 0: + new_cv = _parent_cv(self.cv_stack.pop(), new_cv, self.key, self.flags) + total_chunks >>= 1 + self.cv_stack.append(new_cv) + def update(self, data): + i = 0 + while i < len(data): + if self.chunk_state.len() == CHUNK_LEN: + cv = self.chunk_state.output().chaining_value() + total = self.chunk_state.counter + 1 + self._add_chunk_cv(cv, total) + self.chunk_state = _ChunkState(self.key, total, self.flags) + want = CHUNK_LEN - self.chunk_state.len() + take = min(want, len(data) - i) + self.chunk_state.update(data[i:i+take]) + i += take + def finalize(self, n_bytes=32): + out = self.chunk_state.output() + remaining = list(self.cv_stack) + while remaining: + cv = remaining.pop() + out = _parent_output(cv, out.chaining_value(), self.key, self.flags) + return out.root_output_bytes(n_bytes) + +def blake3_ref(data, out_bytes=32): + h = _Hasher() + h.update(data) + return h.finalize(out_bytes) + + +# ── wire protocol drivers ───────────────────────────────────────────── + +def encode_bsb3(inputs, out_bytes=32): + buf = bytearray() + buf.extend(b"BSB3") + buf.extend(struct.pack("7d} mode={mode}") + ins = gen_inputs(n, mode=mode) + expected = [blake3_ref(x, 32) for x in ins] + got, wall, kernel_ms = run_worker(worker, ins, 32) + total_in = sum(len(x) for x in ins) + mb_s = (total_in / (kernel_ms / 1e3)) / 1e6 if kernel_ms else float('nan') + fail = [] + for i, (a, b) in enumerate(zip(expected, got)): + if a != b: + fail.append((i, a, b)) + if fail: + print(f" FAIL: {len(fail)} mismatches / {n} (kernel {kernel_ms} ms, " + f"wall {wall*1000:.1f} ms, {total_in/1e6:.2f} MB)") + for i, a, b in fail[:max_show_diff]: + print(f" input[{i}] len={len(ins[i])}") + print(f" ref: {a.hex()}") + print(f" got: {b.hex()}") + print(f" in: {ins[i][:32].hex()}...") + return False + print(f" OK kernel={kernel_ms} ms wall={wall*1000:.1f} ms " + f"{total_in/1e6:.2f} MB {mb_s:.1f} MB/s") + return True + + +def main(): + ap = argparse.ArgumentParser() + ap.add_argument("worker", nargs="?", default="./blake3-fanout") + ap.add_argument("--n", type=int, default=None, + help="single n; default = run [32, 1000, 10000, 100000, 1000000]") + ap.add_argument("--mode", default="mixed", + choices=["mixed", "small", "single-chunk", "multi-chunk", "fixed-64"]) + ap.add_argument("--quick", action="store_true", + help="just n=32 mixed + n=1000 multi-chunk") + args = ap.parse_args() + + if not os.path.exists(args.worker): + print(f"worker not found: {args.worker}", file=sys.stderr); sys.exit(2) + + if args.n is not None: + ok = run_case(args.worker, args.n, args.mode) + sys.exit(0 if ok else 1) + if args.quick: + ok = True + ok &= run_case(args.worker, 32, "mixed") + ok &= run_case(args.worker, 1000, "multi-chunk") + sys.exit(0 if ok else 1) + + # full byte-identity sweep across canonical sizes & modes + cases = [ + (32, "mixed"), + (32, "small"), + (32, "single-chunk"), + (32, "multi-chunk"), + (1000, "mixed"), + (10000, "mixed"), + (100000, "fixed-64"), + (1000000, "fixed-64"), + ] + ok = True + for n, mode in cases: + ok &= run_case(args.worker, n, mode) + print("PASS" if ok else "FAIL") + sys.exit(0 if ok else 1) + + +if __name__ == "__main__": + main() diff --git a/examples/cuda-fanout/vendor/blake3-gpu/LICENSE b/examples/cuda-fanout/vendor/blake3-gpu/LICENSE new file mode 100644 index 0000000..30fb3b5 --- /dev/null +++ b/examples/cuda-fanout/vendor/blake3-gpu/LICENSE @@ -0,0 +1,21 @@ +MIT License + +Copyright (c) 2021 Rehan Vipin + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in all +copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +SOFTWARE. diff --git a/examples/cuda-fanout/vendor/blake3-gpu/NOTICE b/examples/cuda-fanout/vendor/blake3-gpu/NOTICE new file mode 100644 index 0000000..c013815 --- /dev/null +++ b/examples/cuda-fanout/vendor/blake3-gpu/NOTICE @@ -0,0 +1,29 @@ +vendor/blake3-gpu/ — BLAKE3 compression primitives derived from +Blaze-3/BLAKE3-gpu. + + Upstream: https://github.com/Blaze-3/BLAKE3-gpu + Commit: f23e8ad889bd472607d39091b9931715d89cc4a8 + License: MIT (see LICENSE in this directory) + Author: Rehan Vipin (2021) + +What we use: + - BLAKE3 IV / message-permutation constants + - g/round/permute/compress primitives translated for __device__ + - leaf-chunk compression (CHUNK_START / CHUNK_END / ROOT flags) + - parent-node compression flag layout + +What we replaced: + - Tree-mode merkle dispatch (dynamic parallelism, recursive kernel + launches, Thrust pinned-vector factory) — our workload is per-item + fan-out where each input is small enough to hash serially inside a + single CUDA thread. We do not need cudaLimitDevRuntimeSyncDepth or + SNICKER staging; we spawn one thread per input and walk the chunk + tree on the device side. SIMT divergence stays bounded by input + length distribution within a warp. + - thrust / iostream / vector includes — pure C kernel uses neither. + +The wrapping `blake3-fanout.cu` binary (host driver + per-input kernel ++ wire-protocol I/O) is AGPLv3, written from scratch in this repo, +under the same protocol as `shake256-fanout.cu` / `cgbn-batch-worker.cu` +/ `secp256k1-batch-mul.cu`. MIT terms above satisfy AGPLv3 inbound +license compatibility (MIT is permissive, no copyleft conflict). diff --git a/examples/cuda-fanout/vendor/blake3-gpu/blake3_device.cuh b/examples/cuda-fanout/vendor/blake3-gpu/blake3_device.cuh new file mode 100644 index 0000000..25a053c --- /dev/null +++ b/examples/cuda-fanout/vendor/blake3-gpu/blake3_device.cuh @@ -0,0 +1,134 @@ +/* blake3_device.cuh — BLAKE3 __device__ primitives. + * + * Derived from Blaze-3/BLAKE3-gpu cuda/blaze3.cuh + cuda/blaze3_cpu.cuh + * (MIT, Rehan Vipin 2021 — see LICENSE in this directory). Stripped + * to the bare compression primitives needed for a single-input, + * single-thread chunk-tree walk. Tree-mode merkle / dynamic-parallelism + * removed — see NOTICE. + * + * Constants and round function track the BLAKE3 spec + * (https://github.com/BLAKE3-team/BLAKE3/blob/master/reference_impl/ + * reference_impl.rs) byte-for-byte. + * + * AGPLv3 on the surrounding wrapper, MIT on this header — see NOTICE. + */ +#ifndef BLAKE3_DEVICE_CUH +#define BLAKE3_DEVICE_CUH + +#include + +namespace blake3_gpu { + +typedef uint8_t u8; +typedef uint32_t u32; +typedef uint64_t u64; + +constexpr u32 OUT_LEN = 32; +constexpr u32 KEY_LEN = 32; +constexpr u32 BLOCK_LEN = 64; +constexpr u32 CHUNK_LEN = 1024; + +constexpr u32 CHUNK_START = 1U << 0; +constexpr u32 CHUNK_END = 1U << 1; +constexpr u32 PARENT = 1U << 2; +constexpr u32 ROOT = 1U << 3; + +__constant__ static const u32 IV[8] = { + 0x6A09E667U, 0xBB67AE85U, 0x3C6EF372U, 0xA54FF53AU, + 0x510E527FU, 0x9B05688CU, 0x1F83D9ABU, 0x5BE0CD19U, +}; + +__constant__ static const int MSG_PERMUTATION[16] = { + 2, 6, 3, 10, 7, 0, 4, 13, + 1, 11, 12, 5, 9, 14, 15, 8, +}; + +__device__ __forceinline__ u32 rotr32(u32 v, int n) { + return (v >> n) | (v << (32 - n)); +} + +__device__ __forceinline__ void g(u32 *st, int a, int b, int c, int d, + u32 mx, u32 my) { + st[a] = st[a] + st[b] + mx; + st[d] = rotr32(st[d] ^ st[a], 16); + st[c] = st[c] + st[d]; + st[b] = rotr32(st[b] ^ st[c], 12); + st[a] = st[a] + st[b] + my; + st[d] = rotr32(st[d] ^ st[a], 8); + st[c] = st[c] + st[d]; + st[b] = rotr32(st[b] ^ st[c], 7); +} + +__device__ __forceinline__ void round_fn(u32 *st, const u32 *m) { + /* columns */ + g(st, 0, 4, 8, 12, m[0], m[1]); + g(st, 1, 5, 9, 13, m[2], m[3]); + g(st, 2, 6, 10, 14, m[4], m[5]); + g(st, 3, 7, 11, 15, m[6], m[7]); + /* diagonals */ + g(st, 0, 5, 10, 15, m[8], m[9]); + g(st, 1, 6, 11, 12, m[10], m[11]); + g(st, 2, 7, 8, 13, m[12], m[13]); + g(st, 3, 4, 9, 14, m[14], m[15]); +} + +__device__ __forceinline__ void permute_msg(u32 *m) { + u32 p[16]; + for (int i = 0; i < 16; i++) p[i] = m[MSG_PERMUTATION[i]]; + for (int i = 0; i < 16; i++) m[i] = p[i]; +} + +/* compress: writes a 16-word state. After 7 rounds and the spec's XOR + * fold, state[0..8] is the chaining value (also the first 32 B of root + * output when the ROOT flag is set). state[8..16] forms the second + * 32 B of root output. Counter / block_len / flags follow the BLAKE3 + * compression-function contract. */ +__device__ __forceinline__ void compress(const u32 *cv, + const u32 *block_words, + u64 counter, + u32 block_len, + u32 flags, + u32 *out16) +{ + u32 st[16]; + st[0] = cv[0]; st[1] = cv[1]; st[2] = cv[2]; st[3] = cv[3]; + st[4] = cv[4]; st[5] = cv[5]; st[6] = cv[6]; st[7] = cv[7]; + st[8] = IV[0]; st[9] = IV[1]; st[10] = IV[2]; st[11] = IV[3]; + st[12] = (u32)counter; + st[13] = (u32)(counter >> 32); + st[14] = block_len; + st[15] = flags; + + u32 m[16]; + for (int i = 0; i < 16; i++) m[i] = block_words[i]; + + round_fn(st, m); permute_msg(m); + round_fn(st, m); permute_msg(m); + round_fn(st, m); permute_msg(m); + round_fn(st, m); permute_msg(m); + round_fn(st, m); permute_msg(m); + round_fn(st, m); permute_msg(m); + round_fn(st, m); /* round 7, no permute after */ + + for (int i = 0; i < 8; i++) { + st[i] ^= st[i + 8]; + st[i + 8] ^= cv[i]; + } + for (int i = 0; i < 16; i++) out16[i] = st[i]; +} + +/* Load up to BLOCK_LEN bytes little-endian into 16 u32 words; tail + * positions zero. */ +__device__ __forceinline__ void words_from_bytes_le(const u8 *bytes, + u32 len, + u32 *words) +{ + for (int i = 0; i < 16; i++) words[i] = 0U; + for (u32 i = 0; i < len; i++) { + words[i >> 2] |= ((u32)bytes[i]) << (8 * (i & 3)); + } +} + +} /* namespace blake3_gpu */ + +#endif /* BLAKE3_DEVICE_CUH */ diff --git a/www/bend.html b/www/bend.html index bd3b6d0..45a2fec 100644 --- a/www/bend.html +++ b/www/bend.html @@ -131,6 +131,12 @@ make gpu-worker LUMBDA=asm # smallest footprint 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] + + cuda-blake3-tree + RTX 3090 + 32.5 GB/s @ 1M × 64 B (kernel 1.97 ms); 20.4 GB/s @ 1k × 1 MB (kernel 51.5 ms); byte-identical to BLAKE3 reference spec across single-chunk + multi-chunk paths + BSB3 binary: out_bytes + n + (u32 len + bytes) per input → BSR3 n + out_bytes + digests + @@ -163,7 +169,7 @@ make gpu-worker LUMBDA=asm # smallest footprint cuda-dilithium-pqsig57.7× keygen+sign+verify vs single CPU threadRTX 3090 TiIACR 2024/1365 cuda-mc-options-pricing25–152×Tesla C1060 / modernGPU Gems Ch.45 cuda-cuFFT-batched-1D8–32× vs MKL; tcFFT 1.1–3.2× vs cuFFTV100 / A100tcFFT - cuda-blake3-tree~5–20× tree modeBlaze-3 CUDABlaze-3 + cuda-blake3-tree32.5 GB/s @ 1M × 64 B; 20.4 GB/s @ 1k × 1 MB on 3090(promoted — see Live)Blaze-3 cuda-bloom-filter-modern~6× CPU; 3.4 B inserts/sB200 / PerlmutterarXiv:2512.15595 cuda-gemm-batched-FP84.8× FP8 vs A100; 716 TFLOPS H100H100 SXMcuBLAS 12.0 cuda-batched-matrix-inverse4.3–16.8× vs MAGMAP100Superfri 2018