diff --git a/docs/bench-log.md b/docs/bench-log.md index b53e8ff51..37d3e2b69 100644 --- a/docs/bench-log.md +++ b/docs/bench-log.md @@ -159,3 +159,27 @@ Attacker speed: delay must only exceed the 2 s publish-or-lose window; margin is Grinding model (3,600 blocks/epoch, advantage uniform 0 to 15%, keep top quartile, one block burned per withheld candidate), gain per epoch in blocks, no delay vs with delay: s=0.1 +0.40 vs 0; s=0.2 +1.66 vs 0; s=0.3 +3.62 (+0.32%, 13.5:1 on burned blocks) vs 0; s=0.4 +6.06 vs 0. Monte Carlo over 2,000,000 epochs agrees to 0.03 blocks. Correctness: NUDUPL, NUCOMP and the Lehmer partial xgcd agree with Cohen 5.4.7 / plain duplication / plain-division xgcd on 15,000 random cases; block prover equals the naive O(T) prover at T = 37, 5,000 and 100,000 in both groups; 216 associativity triples; 3 tamper cases rejected per size. Recommend: class group 1024-bit D from the checkpoint hash, epoch T = 600 x r_ref and era T = 3,600 x r_ref with r_ref the fastest honest single-core rate measured on the devnet (98 M and 588 M on this Mac), fixed at genesis, 20 min lead time for the epoch seed and 2 h for the era draw, 256-bit Fiat-Shamir prime. Open: external review of classgroup.rs against chiavdf, reference core choice, fallback rule for a node without the seed at epoch start, carry D in the proof. + +## 2026-10-03 igneum-pow: Rust crate bit-exact with proto-metal (consensus-engineer) + +Machine: Apple M5 Max, one performance core, rustc 1.99.0 (rustup), release build with LTO. Crate at `igneum-pow/` (seed, generator, memhard, verify, emit; CLI `bench`, `export`, `hash`), standard library only, serde_json as a dev-dependency for the pack tests. +Agreement with the Swift through `proto-cuda/packs/`: program.json instruction by instruction for igneum-genesis, igneum-genesis-mh and igneum-hourly (3 x 64 match); mixer rot/mul/rc match; cache head, last line and FNV-1a 64 `48c4f5bf24166b2e` match; dataset head, `[MASK]` and 64 sampled words match in all three packs; hash vectors 96/96 for igneum-genesis-mh (memory-hard) and 96/96 each for the two closed-form packs. 23 tests, all pass. +Emitted sources: kernel.cu, program.metal, kernel.cl and program.h byte-identical for all three packs, memhard.h and memhard.metal byte-identical for igneum-genesis-mh; `igneum-pow export` then `diff -r` against the packs differs only in the provenance string of vectors.json/vectors.h. +Found: `proto-cuda/packs/igneum-genesis-mh/program.json` is not valid JSON (main.swift line 1291 writes `jhex(cacheLineMask)` inside the `"item"` string). The Rust emitter writes the mask bare and the test normalises that line; fix pending in the Swift. +Cache fill, 256 MiB on one core: 175 to 181 ms in Rust (5 runs) against 184.5 to 190.6 ms Swift and 161.5 ms C++ host reference. +CPU verify per 32-lane warp, avg of 20, 1 GiB dataset: igneum-genesis 0.441 ms (Swift 0.649), /epoch1 0.411 (0.631), /epoch2 0.488 (0.701), igneum-second-seed 0.482 (0.801), igneum-second-seed/epoch1 at 144 loads and 4,608 items 0.579 (1.205). Cold single warps 0.41 to 0.87 ms (Swift 1.16 to 2.11). Closed form 0.002 ms (Swift 0.017). +Reading: Rust is 1.4x to 2.1x faster than the Swift verifier per warp with the same algorithm (register-major lanes, 32-lane interleaved item derivation); 10 ms gate margin about 17x steady, 11x on the worst cold warp. Cache fill is within 5 percent of the Swift and 10 percent slower than clang C++. +API for the fork: `Epoch::memory_hard(seed, day)` once per epoch (fills the cache), then `epoch.hash(nonce)`, `epoch.hash_warp(base)`, `epoch.verify_block(nonce, target)`; `emit::export_pack(&epoch, day, source)` for miner programs. `seed::seed_words_from_bytes` is the boundary for the VDF output. +Not done: no GPU run from Rust; the 256-bit target mapping stays in the fork; the seed is still a string. + +## 3 October 2026, proto-opencl: OpenCL path built and proven without AMD silicon (Apple OpenCL 1.2, pocl, CPU emulator) + +Machine: the same Apple M5 Max. New: `proto-opencl/host.c` (C99, OpenCL 1.2 API), `kernel.cl` in every pack from `--export-pack` (same emitter, OpenCL C dialect; memory-hard core emitted in three dialects), `WAVEFRONT.md`, CPU emulator with a 32- or 64-wide sub-group. The AMD rig has not arrived; no AMD compiler or device has touched this code. +Exchange rule: `sub_group_shuffle_xor` only when the device lists `cl_khr_subgroup_shuffle`, the work-group is exactly 32 and the queried sub-group size for a 32-item work-group is exactly 32; otherwise a `__local` memory exchange with one barrier per exchange (two alternating buffers). Wave64 hardware (GCN, CDNA, RDNA in wave64) therefore takes the local-memory path and the hash never depends on the wave width. +Apple OpenCL 1.2 runtime, Apple M5 Max (40 CUs, OpenCL C 1.2, no sub-group extension, local-memory path): cache check PASS (all 2^26 words, FNV-1a 64 48c4f5bf24166b2e = Mac), dataset self-test PASS, 96/96 vectors standalone and in batch for igneum-genesis-mh; also 96/96 at `--exchange local --group-warps 2` and `--group-warps 4`; closed-form packs igneum-genesis 96/96 and igneum-hourly 96/96. +Apple OpenCL hash rate (wall time; Apple's event timestamps are unusable), pack igneum-genesis-mh, 1 GiB, 5 x 2^24: 45.03 Mhash/s, 18.73 GB/s useful (repeat run 44.58). igneum-genesis 45.17, igneum-hourly 36.33 (128 loads). Metal on the same chip: 45.2. This is Apple's deprecated OpenCL on the M5 Max, NOT an AMD number. +Apple OpenCL sweep (3 batches): 4 MiB 573.7 Mhash/s, 64 MiB 178.9, 256 MiB 94.3, 512 MiB 68.8, 1024 MiB 45.0 (Metal sweep shape reproduced). +pocl 7.2 CPU device (OpenCL 3.0, LLVM 23, Khronos ICD loader, `brew install pocl`, needs SDKROOT): `--exchange auto` and `--exchange subgroup` built with `-cl-std=CL3.0 -D IGNEUM_EXCHANGE=1` and ran the real `sub_group_shuffle_xor` text: cache FNV = Mac, 96/96 PASS. pocl's `clGetKernelSubGroupInfoKHR` returns CL_INVALID_OPERATION, so the new probe kernel (`igneum_probe_subgroup`, reports get_sub_group_size() 32) decided; `--exchange local` also 96/96. +CPU emulator (`proto-opencl/emu`, kernel.cl compiled as C++, 1 GiB dataset built on 256 host threads): 7 configurations all PASS with the identical batch fingerprint f99fb375b3abeaf5 over 2^13 outputs: exchange 0 with work-group 32/sub-group 32, 64/64, 32/64; exchange 1 (sub-group shuffles) with 32/32, 32/64, 64/64 (wave64 carrying two 32-lane units in one shuffle domain), 64/32. +Cross-implementation fingerprint at `--batch-log2 13`, base nonce 0, igneum-genesis-mh: Apple OpenCL f99fb375b3abeaf5, pocl sub-group f99fb375b3abeaf5, pocl local f99fb375b3abeaf5, emulator f99fb375b3abeaf5 (all 7). At 2^24 Apple OpenCL prints 98af644e993239e2 (reference for the AMD run). proto-cuda emulator re-run after the header changes (program.h, vectors.h, memhard.h now C99-safe): PASS. +Not demonstrated: any AMD compile or run, any AMD hash rate, the cost of the local-memory exchange on AMD, whether RDNA compiles igneum_hash as wave32 or wave64. Next: run the seven commands in `proto-opencl/README.md` on the AMD rig and paste the logs. diff --git a/proto-cuda/README.md b/proto-cuda/README.md index 79bc75e77..20a341069 100644 --- a/proto-cuda/README.md +++ b/proto-cuda/README.md @@ -21,7 +21,8 @@ proto-cuda/ CHECKLIST.md Metal/CUDA equivalence, op by op, and what was verified where packs// one program pack per seed, written by proto-metal/igneum-bench --export-pack kernel.cu the program as a CUDA kernel, plus fill kernel and host launch wrappers - program.h seed, day words, dataset size, loads per hash, wrapper declarations + kernel.cl the same program as OpenCL C for proto-opencl (AMD and any other OpenCL device), built at runtime + program.h seed, day words, dataset size, loads per hash, wrapper declarations (C99-safe: proto-opencl/host.c includes it too) vectors.h expected outputs for 3 warps (96 x 64-bit) and dataset self-test values program.json the instruction list and all constants, for any other implementation vectors.json the same vectors as JSON diff --git a/proto-cuda/packs/igneum-genesis-mh/kernel.cl b/proto-cuda/packs/igneum-genesis-mh/kernel.cl new file mode 100644 index 000000000..db3fd0c8b --- /dev/null +++ b/proto-cuda/packs/igneum-genesis-mh/kernel.cl @@ -0,0 +1,278 @@ +// Generated by proto-metal/igneum-bench --export-pack for seed "igneum-genesis". Do not edit by hand. +// OpenCL C twin of the Metal kernel for the same seed (see proto-opencl/README.md, WAVEFRONT.md and program.metal). +// Built from source at runtime by proto-opencl/host.c, which passes these defines: +// IGNEUM_GROUP work-group size of igneum_hash, a multiple of 32 (default 32: one work-group = one 32-lane unit) +// IGNEUM_EXCHANGE 0 = local-memory exchange with a barrier (any device, any wave width; the default) +// 1 = sub_group_shuffle_xor (cl_khr_subgroup_shuffle), only with IGNEUM_GROUP 32 and a sub-group size of exactly 32 +// 2 = intel_sub_group_shuffle_xor (cl_intel_subgroups), same condition +// The verification unit is always 32 lanes. A 64-wide hardware wave (AMD GCN/CDNA, RDNA in wave64) runs two units; +// the exchange masks are 1, 2, 4, 8, 16, so every partner lane lies inside the lane's own aligned run of 32. +#ifndef IGNEUM_GROUP +#define IGNEUM_GROUP 32 +#endif +#ifndef IGNEUM_EXCHANGE +#define IGNEUM_EXCHANGE 0 +#endif +#ifdef __OPENCL_VERSION__ +#define IGNEUM_KERNEL_HASH __kernel __attribute__((reqd_work_group_size(IGNEUM_GROUP, 1, 1))) +#define IGNEUM_LOCAL_WORDS(name, n) __local uint name[n] +#if IGNEUM_EXCHANGE == 1 +#ifdef cl_khr_subgroups +#pragma OPENCL EXTENSION cl_khr_subgroups : enable +#endif +#ifdef cl_khr_subgroup_shuffle +#pragma OPENCL EXTENSION cl_khr_subgroup_shuffle : enable +#endif +#elif IGNEUM_EXCHANGE == 2 +#pragma OPENCL EXTENSION cl_intel_subgroups : enable +#endif +#else +// Not an OpenCL compiler: proto-opencl/emu compiles this file as C++ and supplies the built-ins and these two macros. +#include "emu_opencl.h" +#endif + +#if IGNEUM_EXCHANGE == 1 +#define IGNEUM_SHFL_XOR(dst, a, m) dst = sub_group_shuffle_xor((a), (uint)(m)) +#define IGNEUM_BCAST0(dst, a) dst = sub_group_broadcast((a), 0u) +#elif IGNEUM_EXCHANGE == 2 +#define IGNEUM_SHFL_XOR(dst, a, m) dst = intel_sub_group_shuffle_xor((a), (uint)(m)) +#define IGNEUM_BCAST0(dst, a) dst = sub_group_broadcast((a), 0u) +#else +// Local-memory exchange. Two buffers of IGNEUM_GROUP words alternate (xk counts exchanges), so one barrier per +// exchange is enough: a lane can only overwrite buffer b at exchange k+2 after passing barrier k+1, and every lane +// reaches barrier k+1 only after its read of buffer b at exchange k. The partner lid ^ m stays inside the lane's +// aligned run of 32 because m < 32. Control flow is uniform, so every work-item reaches every barrier. +#define IGNEUM_SHFL_XOR(dst, a, m) { xch[(xk & 1u) * IGNEUM_GROUP + lid] = (a); barrier(CLK_LOCAL_MEM_FENCE); dst = xch[(xk & 1u) * IGNEUM_GROUP + (lid ^ (uint)(m))]; xk += 1u; } +#define IGNEUM_BCAST0(dst, a) { xch[(xk & 1u) * IGNEUM_GROUP + lid] = (a); barrier(CLK_LOCAL_MEM_FENCE); dst = xch[(xk & 1u) * IGNEUM_GROUP + (lid & ~31u)]; xk += 1u; } +#endif + +static inline uint splitmix32(uint x) { + x ^= x >> 16; x *= 0x7feb352du; + x ^= x >> 15; x *= 0x846ca68bu; + x ^= x >> 16; + return x; +} +// n is a literal in 1..31 at every call site. OpenCL rotate() rotates left by n modulo 32. +static inline uint rotl_imm(uint x, uint n) { return rotate(x, n); } +// Right rotation by n modulo 32 as a left rotation by (32 - n) modulo 32; n == 0 gives x. +static inline uint rotr_var(uint x, uint n) { return rotate(x, (0u - n) & 31u); } +static inline uint ds_elem(uint i, uint d0, uint d1) { + uint x = i ^ d0; + x *= 0x9E3779B1u; x ^= x >> 15; + x += d1; + x *= 0x85EBCA77u; x ^= x >> 13; + x *= 0xC2B2AE3Du; x ^= x >> 16; + return x; +} + +// Memory-hard dataset core (MEMHARD.md). Cache: 2^26 words in 2^16 segments of 64 chained ChaCha12 lines. +// Item: 8 rounds of seed-parameterised mixer + one 64-byte cache read, then a final mixer. All parameters are literals. +#define MH_CACHE_LINE_MASK 0x003fffffu +#define MH_SEGMENT_LINES 64u +#define MH_QR(a, b, c, d, r1, r2, r3, r4) { a += b; d ^= a; d = mh_rotl(d, r1); c += d; b ^= c; b = mh_rotl(b, r2); a += b; d ^= a; d = mh_rotl(d, r3); c += d; b ^= c; b = mh_rotl(b, r4); } +static inline uint mh_rotl(uint x, uint n) { return (x << n) | (x >> (32u - n)); } // n in 1..31 at every call site + +// y = ChaCha12 core(x) + x +static inline void mh_chacha_block(const uint* x, uint* y) { + for (uint i = 0u; i < 16u; ++i) y[i] = x[i]; + for (uint r = 0u; r < 6u; ++r) { + MH_QR(y[0], y[4], y[8], y[12], 16u, 12u, 8u, 7u) MH_QR(y[1], y[5], y[9], y[13], 16u, 12u, 8u, 7u) + MH_QR(y[2], y[6], y[10], y[14], 16u, 12u, 8u, 7u) MH_QR(y[3], y[7], y[11], y[15], 16u, 12u, 8u, 7u) + MH_QR(y[0], y[5], y[10], y[15], 16u, 12u, 8u, 7u) MH_QR(y[1], y[6], y[11], y[12], 16u, 12u, 8u, 7u) + MH_QR(y[2], y[7], y[8], y[13], 16u, 12u, 8u, 7u) MH_QR(y[3], y[4], y[9], y[14], 16u, 12u, 8u, 7u) + } + for (uint i = 0u; i < 16u; ++i) y[i] += x[i]; +} + +// One cache segment: 64 chained lines written at cache[seg * 1024]. in_j = prev ^ (sigma || K || seg || j || tag), prev_0 = 0. +static inline void mh_cache_segment(__global uint* cache, uint seg) { + uint prev[16]; uint x[16]; uint y[16]; + for (uint i = 0u; i < 16u; ++i) prev[i] = 0u; + for (uint j = 0u; j < MH_SEGMENT_LINES; ++j) { + x[0] = 0x61707865u ^ prev[0]; x[1] = 0x3320646eu ^ prev[1]; x[2] = 0x79622d32u ^ prev[2]; x[3] = 0x6b206574u ^ prev[3]; + x[4] = 0x3067619fu ^ prev[4]; + x[5] = 0x3c269176u ^ prev[5]; + x[6] = 0x84a03b03u ^ prev[6]; + x[7] = 0xf8c63294u ^ prev[7]; + x[8] = 0xff977c5bu ^ prev[8]; + x[9] = 0xe60def3eu ^ prev[9]; + x[10] = 0x63630141u ^ prev[10]; + x[11] = 0xb8fbcb58u ^ prev[11]; + x[12] = seg ^ prev[12]; x[13] = j ^ prev[13]; x[14] = 0x49676e65u ^ prev[14]; x[15] = 0x756d4d48u ^ prev[15]; + mh_chacha_block(x, y); + __global uint* line = cache + ((seg * MH_SEGMENT_LINES + j) * 16u); + for (uint i = 0u; i < 16u; ++i) { line[i] = y[i]; prev[i] = y[i]; } + } +} + +// M_r: per word (s ^ (RC + rk)) * MUL, then a column round and a diagonal round with the seed-drawn rotations. +static inline void mh_mixer(uint* s, uint rk) { + s[0] = (s[0] ^ (0xbab68293u + rk)) * 0x42146205u; + s[1] = (s[1] ^ (0xcc162340u + rk)) * 0x52cbe0fbu; + s[2] = (s[2] ^ (0x6ce151ccu + rk)) * 0x7ecf4a03u; + s[3] = (s[3] ^ (0xe62b8997u + rk)) * 0x6728907fu; + s[4] = (s[4] ^ (0xc9c80297u + rk)) * 0xd81d9751u; + s[5] = (s[5] ^ (0xf74a1654u + rk)) * 0x132952c3u; + s[6] = (s[6] ^ (0x3d704af5u + rk)) * 0xf60de277u; + s[7] = (s[7] ^ (0x3cf522b7u + rk)) * 0x05358035u; + s[8] = (s[8] ^ (0x2b9cac04u + rk)) * 0xbaf6499du; + s[9] = (s[9] ^ (0xa880ac10u + rk)) * 0xe4db9667u; + s[10] = (s[10] ^ (0x13e5dd1du + rk)) * 0x3e98f45du; + s[11] = (s[11] ^ (0x6fc3e233u + rk)) * 0xd0004eddu; + s[12] = (s[12] ^ (0x2d83eeacu + rk)) * 0x2691630du; + s[13] = (s[13] ^ (0x9006e8bfu + rk)) * 0x9beb3bcfu; + s[14] = (s[14] ^ (0x2c4b5362u + rk)) * 0xab310379u; + s[15] = (s[15] ^ (0x31b49ee2u + rk)) * 0x99cfb423u; + MH_QR(s[0], s[4], s[8], s[12], 20u, 20u, 19u, 4u) MH_QR(s[1], s[5], s[9], s[13], 20u, 20u, 19u, 4u) + MH_QR(s[2], s[6], s[10], s[14], 20u, 20u, 19u, 4u) MH_QR(s[3], s[7], s[11], s[15], 20u, 20u, 19u, 4u) + MH_QR(s[0], s[5], s[10], s[15], 26u, 3u, 3u, 27u) MH_QR(s[1], s[6], s[11], s[12], 26u, 3u, 3u, 27u) + MH_QR(s[2], s[7], s[8], s[13], 26u, 3u, 3u, 27u) MH_QR(s[3], s[4], s[9], s[14], 26u, 3u, 3u, 27u) +} + +// Item t: 16 words. s = (K, t * MUL[i] + RC[i]); 8 rounds of mixer + cache line s[0] & mask; final mixer. +static inline void mh_item(__global const uint* cache, uint t, uint* s) { + s[0] = 0x3067619fu; + s[1] = 0x3c269176u; + s[2] = 0x84a03b03u; + s[3] = 0xf8c63294u; + s[4] = 0xff977c5bu; + s[5] = 0xe60def3eu; + s[6] = 0x63630141u; + s[7] = 0xb8fbcb58u; + s[8] = t * 0x42146205u + 0xbab68293u; + s[9] = t * 0x52cbe0fbu + 0xcc162340u; + s[10] = t * 0x7ecf4a03u + 0x6ce151ccu; + s[11] = t * 0x6728907fu + 0xe62b8997u; + s[12] = t * 0xd81d9751u + 0xc9c80297u; + s[13] = t * 0x132952c3u + 0xf74a1654u; + s[14] = t * 0xf60de277u + 0x3d704af5u; + s[15] = t * 0x05358035u + 0x3cf522b7u; + for (uint r = 0u; r < 8u; ++r) { + mh_mixer(s, 0x9E3779B9u * (r + 1u)); + __global const uint* line = cache + ((s[0] & MH_CACHE_LINE_MASK) * 16u); + for (uint i = 0u; i < 16u; ++i) s[i] ^= line[i]; + } + mh_mixer(s, 0x9E3779B9u * 9u); +} +// dataset[w] without the dataset: derive item w >> 4 and take word w & 15. +static inline uint mh_word(__global const uint* cache, uint w) { uint s[16]; mh_item(cache, w >> 4u, s); return s[w & 15u]; } + +// Memory-hard dataset (MEMHARD.md). One work-item per cache segment; one work-item per 64-byte dataset item. +// The same constants as memhard.h in this pack (one emitter, three dialects). +__kernel void igneum_cache_fill(__global uint* cache, uint nSegments) { + uint seg = (uint)get_global_id(0); + if (seg < nSegments) mh_cache_segment(cache, seg); +} +__kernel void igneum_build(__global uint* ds, __global const uint* cache, uint nItems) { + uint t = (uint)get_global_id(0); + if (t < nItems) { + uint s[16]; + mh_item(cache, t, s); + __global uint* d = ds + ((ulong)t * 16u); + for (uint i = 0u; i < 16u; ++i) d[i] = s[i]; + } +} + +// One hash per work-item. IGNEUM_GROUP is a multiple of 32; lane = lid & 31 and every exchange stays inside the +// lane's own aligned run of 32 work-items, exactly like simd_shuffle_xor inside a 32-wide Metal SIMD group and +// __shfl_xor_sync inside a CUDA warp. Control flow is uniform (no branches at all). +IGNEUM_KERNEL_HASH void igneum_hash(__global const uint* ds, __global ulong* out, uint baseNonce, uint mask) { + uint gid = (uint)get_global_id(0); + uint lid = (uint)get_local_id(0); + uint nonce = baseNonce + gid; + uint r0, r1, r2, r3, r4, r5, r6, r7; +#if IGNEUM_EXCHANGE == 0 + IGNEUM_LOCAL_WORDS(xch, 2 * IGNEUM_GROUP); + uint xk = 0u; +#else + (void)lid; +#endif + { uint x = nonce ^ 0x67a9a7beu; x += 0x9e3779b9u; x = splitmix32(x); r0 = x ^ 0x1a155b25u; } // SEEDW[0], 0x9e3779b9u * 1u, SEEDW[1] + { uint x = nonce ^ 0x1a155b25u; x += 0x3c6ef372u; x = splitmix32(x); r1 = x ^ 0xfddfb732u; } // SEEDW[1], 0x9e3779b9u * 2u, SEEDW[2] + { uint x = nonce ^ 0xfddfb732u; x += 0xdaa66d2bu; x = splitmix32(x); r2 = x ^ 0x4b5af2e8u; } // SEEDW[2], 0x9e3779b9u * 3u, SEEDW[3] + { uint x = nonce ^ 0x4b5af2e8u; x += 0x78dde6e4u; x = splitmix32(x); r3 = x ^ 0xc55caf33u; } // SEEDW[3], 0x9e3779b9u * 4u, SEEDW[4] + { uint x = nonce ^ 0xc55caf33u; x += 0x1715609du; x = splitmix32(x); r4 = x ^ 0xa27c13b7u; } // SEEDW[4], 0x9e3779b9u * 5u, SEEDW[5] + { uint x = nonce ^ 0xa27c13b7u; x += 0xb54cda56u; x = splitmix32(x); r5 = x ^ 0x06628a48u; } // SEEDW[5], 0x9e3779b9u * 6u, SEEDW[6] + { uint x = nonce ^ 0x06628a48u; x += 0x5384540fu; x = splitmix32(x); r6 = x ^ 0x03852469u; } // SEEDW[6], 0x9e3779b9u * 7u, SEEDW[7] + { uint x = nonce ^ 0x03852469u; x += 0xf1bbcdc8u; x = splitmix32(x); r7 = x ^ 0x67a9a7beu; } // SEEDW[7], 0x9e3779b9u * 8u, SEEDW[0] + + for (uint it = 0u; it < 8u; ++it) { + uint sel = r0; + r4 = rotl_imm(r4, 25u); // 0 rotl + r0 = r0 - r5; // 1 sub + r4 = r4 ^ ds[r3 & mask]; // 2 load + r1 = rotl_imm(r1, 1u); // 3 rotl + r2 = r2 + r3 + ((((sel >> 26u) & 1u) != 0u) ? 0x2735a174u : 0x61f0b51cu); // 4 add + r5 = r5 ^ ds[r3 & mask]; // 5 load + r5 = r5 - r7; // 6 sub + r3 = r3 + r4 + ((((sel >> 26u) & 1u) != 0u) ? 0x5a069596u : 0x52f2dbf4u); // 7 add + r0 = r0 ^ r4; // 8 xor + r4 = r4 ^ r0; // 9 xor + r2 = r2 ^ ds[r0 & mask]; // 10 load + { uint t_; IGNEUM_SHFL_XOR(t_, r6, 16u); r4 = r4 ^ t_; } // 11 shfl + r1 = r1 - r5; // 12 sub + r2 = r2 ^ r1; // 13 xor + r4 = r4 ^ ds[r5 & mask]; // 14 load + r2 = r2 ^ ds[r4 & mask]; // 15 load + r3 = r3 ^ ds[r0 & mask]; // 16 load + r4 = r4 ^ r6; // 17 xor + r2 = r4 * r6 + r2; // 18 mad + r6 = rotr_var(r6, r1); // 19 rotr + r3 = r3 ^ r4; // 20 xor + r1 = r3 * r5 + r1; // 21 mad + r7 = mul_hi(r7, r4); // 22 mulhi + r5 = mul_hi(r5, r2); // 23 mulhi + r0 = r0 ^ ds[r6 & mask]; // 24 load + r5 = r5 * r6; // 25 mul + r7 = r7 ^ ds[r1 & mask]; // 26 load + r3 = rotr_var(r3, r1); // 27 rotr + { uint t_; IGNEUM_SHFL_XOR(t_, r1, 8u); r5 = r5 ^ t_; } // 28 shfl + r7 = r7 ^ r5; // 29 xor + r7 = rotl_imm(r7, 23u); // 30 rotl + r2 = r2 - r0; // 31 sub + r7 = r7 ^ r2; // 32 xor + r2 = r2 ^ r6; // 33 xor + r6 = r6 ^ ds[r1 & mask]; // 34 load + r1 = r1 ^ r4; // 35 xor + { uint t_; IGNEUM_SHFL_XOR(t_, r0, 1u); r3 = r3 ^ t_; } // 36 shfl + r2 = r2 * r6; // 37 mul + r5 = r5 + r3 + ((((sel >> 12u) & 1u) != 0u) ? 0xa29f4338u : 0x71f30417u); // 38 add + r7 = r7 ^ r6; // 39 xor + r7 = r7 ^ r3; // 40 xor + r3 = rotr_var(r3, r4); // 41 rotr + r5 = r5 ^ r3; // 42 xor + r3 = rotr_var(r3, r6); // 43 rotr + r1 = r3 * r5 + r1; // 44 mad + r7 = r7 + r3 + ((((sel >> 29u) & 1u) != 0u) ? 0xa907b90bu : 0xc1ae8d3bu); // 45 add + r7 = r7 | r2; // 46 or + { uint t_; IGNEUM_SHFL_XOR(t_, r3, 16u); r2 = r2 ^ t_; } // 47 shfl + r1 = r1 ^ ds[r7 & mask]; // 48 load + r5 = r5 - r1; // 49 sub + r3 = r3 ^ ds[r1 & mask]; // 50 load + r2 = r2 - r3; // 51 sub + { uint t_; IGNEUM_SHFL_XOR(t_, r2, 2u); r6 = r6 ^ t_; } // 52 shfl + r2 = r2 - r0; // 53 sub + r0 = r0 ^ ds[r3 & mask]; // 54 load + { uint t_; IGNEUM_SHFL_XOR(t_, r1, 16u); r2 = r2 ^ t_; } // 55 shfl + r0 = r0 + r6 + ((((sel >> 27u) & 1u) != 0u) ? 0xf4689674u : 0x25955401u); // 56 add + r2 = mul_hi(r2, r0); // 57 mulhi + r4 = mul_hi(r4, r2); // 58 mulhi + r2 = r6 * r7 + r2; // 59 mad + r3 = r3 ^ r1; // 60 xor + r4 = r4 * r2; // 61 mul + r0 = r0 ^ ds[r2 & mask]; // 62 load + r7 = mul_hi(r7, r0); // 63 mulhi + } + uint lo = r0 ^ rotl_imm(r1, 7u) ^ rotl_imm(r2, 14u) ^ rotl_imm(r3, 21u); + uint hi = r4 ^ rotl_imm(r5, 9u) ^ rotl_imm(r6, 18u) ^ rotl_imm(r7, 27u); + out[gid] = ((ulong)hi << 32) | (ulong)lo; +} + +#if IGNEUM_EXCHANGE != 0 +// Reports the sub-group size this device uses for a work-group of IGNEUM_GROUP items. host.c runs it only when the +// per-kernel query (clGetKernelSubGroupInfoKHR on igneum_hash) is unavailable; that query is preferred because a +// compiler may pick a different wave width per kernel (RDNA: wave32 or wave64). See WAVEFRONT.md. +IGNEUM_KERNEL_HASH void igneum_probe_subgroup(__global uint* out) { + if (get_local_id(0) == 0u) { out[0] = get_sub_group_size(); out[1] = get_num_sub_groups(); } +} +#endif diff --git a/proto-cuda/packs/igneum-genesis-mh/memhard.h b/proto-cuda/packs/igneum-genesis-mh/memhard.h index 41aca0de6..6945a281b 100644 --- a/proto-cuda/packs/igneum-genesis-mh/memhard.h +++ b/proto-cuda/packs/igneum-genesis-mh/memhard.h @@ -1,10 +1,17 @@ // Generated by proto-metal/igneum-bench --export-pack for seed "igneum-genesis". Do not edit by hand. // Memory-hard dataset core, the same text that the Mac's Metal kernels and CPU verifier were checked against. -// Included by kernel.cu (device) and host.cu (host reference). See proto-metal/MEMHARD.md for the construction. +// Included by kernel.cu (device), host.cu (host reference) and proto-opencl/host.c (C99 host reference). +// See proto-metal/MEMHARD.md for the construction. kernel.cl carries the same text in OpenCL C. #pragma once +#ifdef __cplusplus #include +#else +#include +#endif #if defined(__CUDACC__) #define IGNEUM_HD __host__ __device__ __forceinline__ +#elif defined(_MSC_VER) && !defined(__cplusplus) +#define IGNEUM_HD static __inline #else #define IGNEUM_HD static inline #endif diff --git a/proto-cuda/packs/igneum-genesis-mh/program.h b/proto-cuda/packs/igneum-genesis-mh/program.h index 86c5e91c5..a2acba1b4 100644 --- a/proto-cuda/packs/igneum-genesis-mh/program.h +++ b/proto-cuda/packs/igneum-genesis-mh/program.h @@ -1,8 +1,15 @@ // Generated by proto-metal/igneum-bench --export-pack for seed "igneum-genesis". Do not edit by hand. // Program metadata for host.cu plus the launch wrappers defined in kernel.cu. +// Also included by proto-opencl/host.c (C99), which defines IGNEUM_NO_CUDA first and reads only the macros. #pragma once -#include +#ifdef __cplusplus #include +#else +#include +#endif +#ifndef IGNEUM_NO_CUDA +#include +#endif #define IGNEUM_SEED_STRING "igneum-genesis" #define IGNEUM_DAY_STRING "2026-10-03" @@ -29,9 +36,11 @@ #define IGNEUM_MIX_MUL_INIT { 0x42146205u, 0x52cbe0fbu, 0x7ecf4a03u, 0x6728907fu, 0xd81d9751u, 0x132952c3u, 0xf60de277u, 0x05358035u, 0xbaf6499du, 0xe4db9667u, 0x3e98f45du, 0xd0004eddu, 0x2691630du, 0x9beb3bcfu, 0xab310379u, 0x99cfb423u } #define IGNEUM_MIX_RC_INIT { 0xbab68293u, 0xcc162340u, 0x6ce151ccu, 0xe62b8997u, 0xc9c80297u, 0xf74a1654u, 0x3d704af5u, 0x3cf522b7u, 0x2b9cac04u, 0xa880ac10u, 0x13e5dd1du, 0x6fc3e233u, 0x2d83eeacu, 0x9006e8bfu, 0x2c4b5362u, 0x31b49ee2u } +#ifndef IGNEUM_NO_CUDA // Defined in kernel.cu. All launch on the default stream and return cudaGetLastError(). cudaError_t igneum_launch_cache_fill(uint32_t* cache, uint32_t nSegments); cudaError_t igneum_launch_build(uint32_t* ds, const uint32_t* cache, uint32_t nItems); cudaError_t igneum_launch_hash(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask, uint32_t nonces, uint32_t blockWarps); cudaError_t igneum_hash_info(int* numRegs, int* blocksPerSM, uint32_t blockWarps); +#endif diff --git a/proto-cuda/packs/igneum-genesis-mh/vectors.h b/proto-cuda/packs/igneum-genesis-mh/vectors.h index 2c97eb016..08245b5ff 100644 --- a/proto-cuda/packs/igneum-genesis-mh/vectors.h +++ b/proto-cuda/packs/igneum-genesis-mh/vectors.h @@ -1,7 +1,11 @@ // Generated by proto-metal/igneum-bench --export-pack for seed "igneum-genesis". Do not edit by hand. // Expected outputs: proto-metal CPU interpreter (cpuWarp, memory-hard dataset) on Apple M5 Max; Metal GPU cross-check PASS 3/3 warps #pragma once +#ifdef __cplusplus #include +#else +#include +#endif #define IGNEUM_VEC_WARPS 3 static const uint32_t IGNEUM_VEC_BASE[IGNEUM_VEC_WARPS] = { 0u, 4096u, 1000000u }; diff --git a/proto-cuda/packs/igneum-genesis/kernel.cl b/proto-cuda/packs/igneum-genesis/kernel.cl new file mode 100644 index 000000000..1a1a36457 --- /dev/null +++ b/proto-cuda/packs/igneum-genesis/kernel.cl @@ -0,0 +1,176 @@ +// Generated by proto-metal/igneum-bench --export-pack for seed "igneum-genesis". Do not edit by hand. +// OpenCL C twin of the Metal kernel for the same seed (see proto-opencl/README.md, WAVEFRONT.md and program.metal). +// Built from source at runtime by proto-opencl/host.c, which passes these defines: +// IGNEUM_GROUP work-group size of igneum_hash, a multiple of 32 (default 32: one work-group = one 32-lane unit) +// IGNEUM_EXCHANGE 0 = local-memory exchange with a barrier (any device, any wave width; the default) +// 1 = sub_group_shuffle_xor (cl_khr_subgroup_shuffle), only with IGNEUM_GROUP 32 and a sub-group size of exactly 32 +// 2 = intel_sub_group_shuffle_xor (cl_intel_subgroups), same condition +// The verification unit is always 32 lanes. A 64-wide hardware wave (AMD GCN/CDNA, RDNA in wave64) runs two units; +// the exchange masks are 1, 2, 4, 8, 16, so every partner lane lies inside the lane's own aligned run of 32. +#ifndef IGNEUM_GROUP +#define IGNEUM_GROUP 32 +#endif +#ifndef IGNEUM_EXCHANGE +#define IGNEUM_EXCHANGE 0 +#endif +#ifdef __OPENCL_VERSION__ +#define IGNEUM_KERNEL_HASH __kernel __attribute__((reqd_work_group_size(IGNEUM_GROUP, 1, 1))) +#define IGNEUM_LOCAL_WORDS(name, n) __local uint name[n] +#if IGNEUM_EXCHANGE == 1 +#ifdef cl_khr_subgroups +#pragma OPENCL EXTENSION cl_khr_subgroups : enable +#endif +#ifdef cl_khr_subgroup_shuffle +#pragma OPENCL EXTENSION cl_khr_subgroup_shuffle : enable +#endif +#elif IGNEUM_EXCHANGE == 2 +#pragma OPENCL EXTENSION cl_intel_subgroups : enable +#endif +#else +// Not an OpenCL compiler: proto-opencl/emu compiles this file as C++ and supplies the built-ins and these two macros. +#include "emu_opencl.h" +#endif + +#if IGNEUM_EXCHANGE == 1 +#define IGNEUM_SHFL_XOR(dst, a, m) dst = sub_group_shuffle_xor((a), (uint)(m)) +#define IGNEUM_BCAST0(dst, a) dst = sub_group_broadcast((a), 0u) +#elif IGNEUM_EXCHANGE == 2 +#define IGNEUM_SHFL_XOR(dst, a, m) dst = intel_sub_group_shuffle_xor((a), (uint)(m)) +#define IGNEUM_BCAST0(dst, a) dst = sub_group_broadcast((a), 0u) +#else +// Local-memory exchange. Two buffers of IGNEUM_GROUP words alternate (xk counts exchanges), so one barrier per +// exchange is enough: a lane can only overwrite buffer b at exchange k+2 after passing barrier k+1, and every lane +// reaches barrier k+1 only after its read of buffer b at exchange k. The partner lid ^ m stays inside the lane's +// aligned run of 32 because m < 32. Control flow is uniform, so every work-item reaches every barrier. +#define IGNEUM_SHFL_XOR(dst, a, m) { xch[(xk & 1u) * IGNEUM_GROUP + lid] = (a); barrier(CLK_LOCAL_MEM_FENCE); dst = xch[(xk & 1u) * IGNEUM_GROUP + (lid ^ (uint)(m))]; xk += 1u; } +#define IGNEUM_BCAST0(dst, a) { xch[(xk & 1u) * IGNEUM_GROUP + lid] = (a); barrier(CLK_LOCAL_MEM_FENCE); dst = xch[(xk & 1u) * IGNEUM_GROUP + (lid & ~31u)]; xk += 1u; } +#endif + +static inline uint splitmix32(uint x) { + x ^= x >> 16; x *= 0x7feb352du; + x ^= x >> 15; x *= 0x846ca68bu; + x ^= x >> 16; + return x; +} +// n is a literal in 1..31 at every call site. OpenCL rotate() rotates left by n modulo 32. +static inline uint rotl_imm(uint x, uint n) { return rotate(x, n); } +// Right rotation by n modulo 32 as a left rotation by (32 - n) modulo 32; n == 0 gives x. +static inline uint rotr_var(uint x, uint n) { return rotate(x, (0u - n) & 31u); } +static inline uint ds_elem(uint i, uint d0, uint d1) { + uint x = i ^ d0; + x *= 0x9E3779B1u; x ^= x >> 15; + x += d1; + x *= 0x85EBCA77u; x ^= x >> 13; + x *= 0xC2B2AE3Du; x ^= x >> 16; + return x; +} + +// dataset[i] = ds_elem(i, d0, d1) for i < n. Same closed form as the Metal igneum_fill kernel. +__kernel void igneum_fill(__global uint* ds, uint n, uint d0, uint d1) { + uint i = (uint)get_global_id(0); + if (i < n) ds[i] = ds_elem(i, d0, d1); +} + +// One hash per work-item. IGNEUM_GROUP is a multiple of 32; lane = lid & 31 and every exchange stays inside the +// lane's own aligned run of 32 work-items, exactly like simd_shuffle_xor inside a 32-wide Metal SIMD group and +// __shfl_xor_sync inside a CUDA warp. Control flow is uniform (no branches at all). +IGNEUM_KERNEL_HASH void igneum_hash(__global const uint* ds, __global ulong* out, uint baseNonce, uint mask) { + uint gid = (uint)get_global_id(0); + uint lid = (uint)get_local_id(0); + uint nonce = baseNonce + gid; + uint r0, r1, r2, r3, r4, r5, r6, r7; +#if IGNEUM_EXCHANGE == 0 + IGNEUM_LOCAL_WORDS(xch, 2 * IGNEUM_GROUP); + uint xk = 0u; +#else + (void)lid; +#endif + { uint x = nonce ^ 0x67a9a7beu; x += 0x9e3779b9u; x = splitmix32(x); r0 = x ^ 0x1a155b25u; } // SEEDW[0], 0x9e3779b9u * 1u, SEEDW[1] + { uint x = nonce ^ 0x1a155b25u; x += 0x3c6ef372u; x = splitmix32(x); r1 = x ^ 0xfddfb732u; } // SEEDW[1], 0x9e3779b9u * 2u, SEEDW[2] + { uint x = nonce ^ 0xfddfb732u; x += 0xdaa66d2bu; x = splitmix32(x); r2 = x ^ 0x4b5af2e8u; } // SEEDW[2], 0x9e3779b9u * 3u, SEEDW[3] + { uint x = nonce ^ 0x4b5af2e8u; x += 0x78dde6e4u; x = splitmix32(x); r3 = x ^ 0xc55caf33u; } // SEEDW[3], 0x9e3779b9u * 4u, SEEDW[4] + { uint x = nonce ^ 0xc55caf33u; x += 0x1715609du; x = splitmix32(x); r4 = x ^ 0xa27c13b7u; } // SEEDW[4], 0x9e3779b9u * 5u, SEEDW[5] + { uint x = nonce ^ 0xa27c13b7u; x += 0xb54cda56u; x = splitmix32(x); r5 = x ^ 0x06628a48u; } // SEEDW[5], 0x9e3779b9u * 6u, SEEDW[6] + { uint x = nonce ^ 0x06628a48u; x += 0x5384540fu; x = splitmix32(x); r6 = x ^ 0x03852469u; } // SEEDW[6], 0x9e3779b9u * 7u, SEEDW[7] + { uint x = nonce ^ 0x03852469u; x += 0xf1bbcdc8u; x = splitmix32(x); r7 = x ^ 0x67a9a7beu; } // SEEDW[7], 0x9e3779b9u * 8u, SEEDW[0] + + for (uint it = 0u; it < 8u; ++it) { + uint sel = r0; + r4 = rotl_imm(r4, 25u); // 0 rotl + r0 = r0 - r5; // 1 sub + r4 = r4 ^ ds[r3 & mask]; // 2 load + r1 = rotl_imm(r1, 1u); // 3 rotl + r2 = r2 + r3 + ((((sel >> 26u) & 1u) != 0u) ? 0x2735a174u : 0x61f0b51cu); // 4 add + r5 = r5 ^ ds[r3 & mask]; // 5 load + r5 = r5 - r7; // 6 sub + r3 = r3 + r4 + ((((sel >> 26u) & 1u) != 0u) ? 0x5a069596u : 0x52f2dbf4u); // 7 add + r0 = r0 ^ r4; // 8 xor + r4 = r4 ^ r0; // 9 xor + r2 = r2 ^ ds[r0 & mask]; // 10 load + { uint t_; IGNEUM_SHFL_XOR(t_, r6, 16u); r4 = r4 ^ t_; } // 11 shfl + r1 = r1 - r5; // 12 sub + r2 = r2 ^ r1; // 13 xor + r4 = r4 ^ ds[r5 & mask]; // 14 load + r2 = r2 ^ ds[r4 & mask]; // 15 load + r3 = r3 ^ ds[r0 & mask]; // 16 load + r4 = r4 ^ r6; // 17 xor + r2 = r4 * r6 + r2; // 18 mad + r6 = rotr_var(r6, r1); // 19 rotr + r3 = r3 ^ r4; // 20 xor + r1 = r3 * r5 + r1; // 21 mad + r7 = mul_hi(r7, r4); // 22 mulhi + r5 = mul_hi(r5, r2); // 23 mulhi + r0 = r0 ^ ds[r6 & mask]; // 24 load + r5 = r5 * r6; // 25 mul + r7 = r7 ^ ds[r1 & mask]; // 26 load + r3 = rotr_var(r3, r1); // 27 rotr + { uint t_; IGNEUM_SHFL_XOR(t_, r1, 8u); r5 = r5 ^ t_; } // 28 shfl + r7 = r7 ^ r5; // 29 xor + r7 = rotl_imm(r7, 23u); // 30 rotl + r2 = r2 - r0; // 31 sub + r7 = r7 ^ r2; // 32 xor + r2 = r2 ^ r6; // 33 xor + r6 = r6 ^ ds[r1 & mask]; // 34 load + r1 = r1 ^ r4; // 35 xor + { uint t_; IGNEUM_SHFL_XOR(t_, r0, 1u); r3 = r3 ^ t_; } // 36 shfl + r2 = r2 * r6; // 37 mul + r5 = r5 + r3 + ((((sel >> 12u) & 1u) != 0u) ? 0xa29f4338u : 0x71f30417u); // 38 add + r7 = r7 ^ r6; // 39 xor + r7 = r7 ^ r3; // 40 xor + r3 = rotr_var(r3, r4); // 41 rotr + r5 = r5 ^ r3; // 42 xor + r3 = rotr_var(r3, r6); // 43 rotr + r1 = r3 * r5 + r1; // 44 mad + r7 = r7 + r3 + ((((sel >> 29u) & 1u) != 0u) ? 0xa907b90bu : 0xc1ae8d3bu); // 45 add + r7 = r7 | r2; // 46 or + { uint t_; IGNEUM_SHFL_XOR(t_, r3, 16u); r2 = r2 ^ t_; } // 47 shfl + r1 = r1 ^ ds[r7 & mask]; // 48 load + r5 = r5 - r1; // 49 sub + r3 = r3 ^ ds[r1 & mask]; // 50 load + r2 = r2 - r3; // 51 sub + { uint t_; IGNEUM_SHFL_XOR(t_, r2, 2u); r6 = r6 ^ t_; } // 52 shfl + r2 = r2 - r0; // 53 sub + r0 = r0 ^ ds[r3 & mask]; // 54 load + { uint t_; IGNEUM_SHFL_XOR(t_, r1, 16u); r2 = r2 ^ t_; } // 55 shfl + r0 = r0 + r6 + ((((sel >> 27u) & 1u) != 0u) ? 0xf4689674u : 0x25955401u); // 56 add + r2 = mul_hi(r2, r0); // 57 mulhi + r4 = mul_hi(r4, r2); // 58 mulhi + r2 = r6 * r7 + r2; // 59 mad + r3 = r3 ^ r1; // 60 xor + r4 = r4 * r2; // 61 mul + r0 = r0 ^ ds[r2 & mask]; // 62 load + r7 = mul_hi(r7, r0); // 63 mulhi + } + uint lo = r0 ^ rotl_imm(r1, 7u) ^ rotl_imm(r2, 14u) ^ rotl_imm(r3, 21u); + uint hi = r4 ^ rotl_imm(r5, 9u) ^ rotl_imm(r6, 18u) ^ rotl_imm(r7, 27u); + out[gid] = ((ulong)hi << 32) | (ulong)lo; +} + +#if IGNEUM_EXCHANGE != 0 +// Reports the sub-group size this device uses for a work-group of IGNEUM_GROUP items. host.c runs it only when the +// per-kernel query (clGetKernelSubGroupInfoKHR on igneum_hash) is unavailable; that query is preferred because a +// compiler may pick a different wave width per kernel (RDNA: wave32 or wave64). See WAVEFRONT.md. +IGNEUM_KERNEL_HASH void igneum_probe_subgroup(__global uint* out) { + if (get_local_id(0) == 0u) { out[0] = get_sub_group_size(); out[1] = get_num_sub_groups(); } +} +#endif diff --git a/proto-cuda/packs/igneum-genesis/program.h b/proto-cuda/packs/igneum-genesis/program.h index 0a4bf2964..172a5917a 100644 --- a/proto-cuda/packs/igneum-genesis/program.h +++ b/proto-cuda/packs/igneum-genesis/program.h @@ -1,8 +1,15 @@ // Generated by proto-metal/igneum-bench --export-pack for seed "igneum-genesis". Do not edit by hand. // Program metadata for host.cu plus the launch wrappers defined in kernel.cu. +// Also included by proto-opencl/host.c (C99), which defines IGNEUM_NO_CUDA first and reads only the macros. #pragma once -#include +#ifdef __cplusplus #include +#else +#include +#endif +#ifndef IGNEUM_NO_CUDA +#include +#endif #define IGNEUM_SEED_STRING "igneum-genesis" #define IGNEUM_DAY_STRING "2026-10-03" @@ -14,12 +21,16 @@ #define IGNEUM_ITERATIONS 8 #define IGNEUM_INSTR_COUNT 64 #define IGNEUM_LOADS_PER_HASH 104 +#define IGNEUM_WIDE_LOADS_PER_HASH 0 #define IGNEUM_OP_MIX "load=13 xor=13 sub=7 shfl=6 add=5 mulhi=5 mad=4 rotr=4 mul=3 rotl=3 or=1" +// 0 = closed-form dataset (ds_elem), 1 = memory-hard cache construction (MEMHARD.md, memhard.h) +#define IGNEUM_DATASET_MODE 0 #define IGNEUM_SEEDW_INIT { 0x67a9a7beu, 0x1a155b25u, 0xfddfb732u, 0x4b5af2e8u, 0xc55caf33u, 0xa27c13b7u, 0x06628a48u, 0x03852469u } - +#ifndef IGNEUM_NO_CUDA // Defined in kernel.cu. Both launch on the default stream and return cudaGetLastError(). cudaError_t igneum_launch_fill(uint32_t* ds, uint32_t nWords, uint32_t d0, uint32_t d1); cudaError_t igneum_launch_hash(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask, uint32_t nonces, uint32_t blockWarps); cudaError_t igneum_hash_info(int* numRegs, int* blocksPerSM, uint32_t blockWarps); +#endif diff --git a/proto-cuda/packs/igneum-genesis/program.json b/proto-cuda/packs/igneum-genesis/program.json index c38266f92..c300524b7 100644 --- a/proto-cuda/packs/igneum-genesis/program.json +++ b/proto-cuda/packs/igneum-genesis/program.json @@ -1,5 +1,6 @@ { - "format": "igneum-program-pack-1", + "format": "igneum-program-pack-2", + "dataset_mode": "closed-form", "seed": "igneum-genesis", "seed_words": ["0x67a9a7be", "0x1a155b25", "0xfddfb732", "0x4b5af2e8", "0xc55caf33", "0xa27c13b7", "0x06628a48", "0x03852469"], "seed_derivation": "FNV-1a 64 over UTF-8 of seed, basis ^ (salt * 0x9E3779B97F4A7C15) for salt 0..3, then h ^= h>>33; h *= 0xff51afd7ed558ccd; h ^= h>>33; words[2*salt] = low 32, words[2*salt+1] = high 32", @@ -24,7 +25,8 @@ "rotr": "dst = rotr(dst, src & 31)", "mad": "dst = src * src2 + dst", "shfl": "dst = dst ^ (src of lane (lane ^ mask)), mask in {1,2,4,8,16}, within the 32-lane warp", - "load": "dst = dst ^ dataset[src & dataset.mask]" + "load": "dst = dst ^ dataset[src & dataset.mask]", + "wload": "base = (src of lane 0 & dataset.mask) & ~31; dst = dst ^ dataset[base + lane] (warp-coalesced 128-byte load, lever b, only when --wide-frac > 0)" }, "dataset": { "log2_words": 28, @@ -34,6 +36,7 @@ "day_words_from": "day/2026-10-03", "d0": "0x3067619f", "d1": "0x3c269176", + "mode": "closed-form", "formula": "x = i ^ d0; x *= 0x9E3779B1; x ^= x>>15; x += d1; x *= 0x85EBCA77; x ^= x>>13; x *= 0xC2B2AE3D; x ^= x>>16 (all mod 2^32)" }, "instructions": [ diff --git a/proto-cuda/packs/igneum-genesis/vectors.h b/proto-cuda/packs/igneum-genesis/vectors.h index 1634ffdd4..1ef33d3c9 100644 --- a/proto-cuda/packs/igneum-genesis/vectors.h +++ b/proto-cuda/packs/igneum-genesis/vectors.h @@ -1,7 +1,11 @@ // Generated by proto-metal/igneum-bench --export-pack for seed "igneum-genesis". Do not edit by hand. -// Expected outputs: proto-metal CPU interpreter (cpuWarp) on Apple M5 Max; Metal GPU cross-check PASS 3/3 warps +// Expected outputs: proto-metal CPU interpreter (cpuWarp, closed-form dataset) on Apple M5 Max; Metal GPU cross-check PASS 3/3 warps #pragma once +#ifdef __cplusplus #include +#else +#include +#endif #define IGNEUM_VEC_WARPS 3 static const uint32_t IGNEUM_VEC_BASE[IGNEUM_VEC_WARPS] = { 0u, 4096u, 1000000u }; @@ -33,3 +37,11 @@ static const uint32_t IGNEUM_DS_HEAD[16] = { }; static const uint32_t IGNEUM_DS_LAST_INDEX = 268435455u; static const uint32_t IGNEUM_DS_LAST = 0xf78c84a4u; +// 64 sampled dataset words (index, value) computed on the Mac. +#define IGNEUM_DS_SAMPLES 64 +static const uint32_t IGNEUM_DS_SAMPLE_INDEX[IGNEUM_DS_SAMPLES] = { + 59471966u, 217795994u, 208353206u, 42483309u, 172547758u, 148076330u, 183853158u, 214389424u, 267488061u, 169781097u, 184093494u, 153880993u, 84977930u, 46426879u, 3093825u, 225364072u, 44593546u, 260713159u, 168250303u, 52384140u, 223401610u, 45554030u, 95410555u, 175039924u, 79171087u, 267580473u, 24168642u, 37981670u, 171551130u, 195559979u, 204611762u, 140997658u, 138925853u, 86637313u, 20736778u, 219665210u, 160430336u, 264654675u, 8013395u, 228945585u, 213884386u, 104419827u, 44185464u, 142737231u, 99284897u, 132475900u, 61861762u, 132056166u, 262388043u, 91878046u, 117353561u, 124768597u, 71352993u, 190698941u, 46055428u, 55281366u, 165145231u, 106810753u, 171985651u, 232085256u, 159510492u, 40072060u, 209107596u, 39023794u +}; +static const uint32_t IGNEUM_DS_SAMPLE_VALUE[IGNEUM_DS_SAMPLES] = { + 0x3ef2f18du, 0xed632a6bu, 0x2daad897u, 0x1201639fu, 0x8e79ca54u, 0x488d4f3du, 0x5206fa5cu, 0xf60c5f15u, 0x5af790f7u, 0xdb48b2cau, 0x2b55a0e4u, 0xc8c5100fu, 0xe4afb4a4u, 0x35a42586u, 0xff4cb6d1u, 0x1d6cac9fu, 0x889230c1u, 0x95316e2au, 0x61b099e9u, 0x97b89bb2u, 0x05171be5u, 0x7f0157abu, 0x13825368u, 0x562ab1bfu, 0x6e3b06b8u, 0xc9b51434u, 0xc994b628u, 0x26643aa8u, 0x56ae558cu, 0xdc8fdaaau, 0x2ed21c3fu, 0x3a29b91bu, 0xfd98ca1eu, 0x44c7dd8au, 0xaab155c8u, 0x0db02165u, 0x637c813au, 0x5648617eu, 0x3e163337u, 0x64bf0ad7u, 0x43381384u, 0x824dca58u, 0xab7bf62eu, 0x10880075u, 0x292498dau, 0x925fea4du, 0x606466a0u, 0x37923fc4u, 0x3544915bu, 0x742b205du, 0x3228e7adu, 0x1ce5ec92u, 0xd7857a5cu, 0x200a560bu, 0xc29076c7u, 0x886a1a06u, 0x1f415038u, 0x60fa6028u, 0x1fb1d0f9u, 0xee203b87u, 0xa12f7ea5u, 0x6ecd91dcu, 0x260e00b4u, 0xef4a4067u +}; diff --git a/proto-cuda/packs/igneum-genesis/vectors.json b/proto-cuda/packs/igneum-genesis/vectors.json index 15350bfd1..90a3660b3 100644 --- a/proto-cuda/packs/igneum-genesis/vectors.json +++ b/proto-cuda/packs/igneum-genesis/vectors.json @@ -1,10 +1,11 @@ { "seed": "igneum-genesis", "day": "2026-10-03", + "dataset_mode": "closed-form", "dataset_log2_words": 28, "mask": "0x0fffffff", "lanes": 32, - "source": "proto-metal CPU interpreter (cpuWarp) on Apple M5 Max; Metal GPU cross-check PASS 3/3 warps", + "source": "proto-metal CPU interpreter (cpuWarp, closed-form dataset) on Apple M5 Max; Metal GPU cross-check PASS 3/3 warps", "warps": [ {"base_nonce": 0, "expected": [ "0x2941e93c76cb1910", "0xc0099c34df8280ae", "0xe61da8a237797181", "0xfb944794eaa8ea1d", "0x0093934e41befb07", "0xd4ba2031c0385915", "0x75a8a43902cee166", "0xb4ce2d6829aaa27c", @@ -27,5 +28,6 @@ ], "dataset_head": ["0x82174c0f", "0x577bdb9c", "0x111053bf", "0x2bb85514", "0xd83e190f", "0x8ccd9427", "0x3f69d5e4", "0xf3bbe1dc", "0xb3aa7e90", "0xc33f5f73", "0xb83a2b10", "0xc4d7c8ff", "0xefa6d1a8", "0x7029b116", "0x5e48fab0", "0x0ef66e40"], "dataset_last_index": 268435455, - "dataset_last": "0xf78c84a4" + "dataset_last": "0xf78c84a4", + "dataset_samples": [{"index": 59471966, "value": "0x3ef2f18d"}, {"index": 217795994, "value": "0xed632a6b"}, {"index": 208353206, "value": "0x2daad897"}, {"index": 42483309, "value": "0x1201639f"}, {"index": 172547758, "value": "0x8e79ca54"}, {"index": 148076330, "value": "0x488d4f3d"}, {"index": 183853158, "value": "0x5206fa5c"}, {"index": 214389424, "value": "0xf60c5f15"}, {"index": 267488061, "value": "0x5af790f7"}, {"index": 169781097, "value": "0xdb48b2ca"}, {"index": 184093494, "value": "0x2b55a0e4"}, {"index": 153880993, "value": "0xc8c5100f"}, {"index": 84977930, "value": "0xe4afb4a4"}, {"index": 46426879, "value": "0x35a42586"}, {"index": 3093825, "value": "0xff4cb6d1"}, {"index": 225364072, "value": "0x1d6cac9f"}, {"index": 44593546, "value": "0x889230c1"}, {"index": 260713159, "value": "0x95316e2a"}, {"index": 168250303, "value": "0x61b099e9"}, {"index": 52384140, "value": "0x97b89bb2"}, {"index": 223401610, "value": "0x05171be5"}, {"index": 45554030, "value": "0x7f0157ab"}, {"index": 95410555, "value": "0x13825368"}, {"index": 175039924, "value": "0x562ab1bf"}, {"index": 79171087, "value": "0x6e3b06b8"}, {"index": 267580473, "value": "0xc9b51434"}, {"index": 24168642, "value": "0xc994b628"}, {"index": 37981670, "value": "0x26643aa8"}, {"index": 171551130, "value": "0x56ae558c"}, {"index": 195559979, "value": "0xdc8fdaaa"}, {"index": 204611762, "value": "0x2ed21c3f"}, {"index": 140997658, "value": "0x3a29b91b"}, {"index": 138925853, "value": "0xfd98ca1e"}, {"index": 86637313, "value": "0x44c7dd8a"}, {"index": 20736778, "value": "0xaab155c8"}, {"index": 219665210, "value": "0x0db02165"}, {"index": 160430336, "value": "0x637c813a"}, {"index": 264654675, "value": "0x5648617e"}, {"index": 8013395, "value": "0x3e163337"}, {"index": 228945585, "value": "0x64bf0ad7"}, {"index": 213884386, "value": "0x43381384"}, {"index": 104419827, "value": "0x824dca58"}, {"index": 44185464, "value": "0xab7bf62e"}, {"index": 142737231, "value": "0x10880075"}, {"index": 99284897, "value": "0x292498da"}, {"index": 132475900, "value": "0x925fea4d"}, {"index": 61861762, "value": "0x606466a0"}, {"index": 132056166, "value": "0x37923fc4"}, {"index": 262388043, "value": "0x3544915b"}, {"index": 91878046, "value": "0x742b205d"}, {"index": 117353561, "value": "0x3228e7ad"}, {"index": 124768597, "value": "0x1ce5ec92"}, {"index": 71352993, "value": "0xd7857a5c"}, {"index": 190698941, "value": "0x200a560b"}, {"index": 46055428, "value": "0xc29076c7"}, {"index": 55281366, "value": "0x886a1a06"}, {"index": 165145231, "value": "0x1f415038"}, {"index": 106810753, "value": "0x60fa6028"}, {"index": 171985651, "value": "0x1fb1d0f9"}, {"index": 232085256, "value": "0xee203b87"}, {"index": 159510492, "value": "0xa12f7ea5"}, {"index": 40072060, "value": "0x6ecd91dc"}, {"index": 209107596, "value": "0x260e00b4"}, {"index": 39023794, "value": "0xef4a4067"}] } diff --git a/proto-cuda/packs/igneum-hourly/kernel.cl b/proto-cuda/packs/igneum-hourly/kernel.cl new file mode 100644 index 000000000..d22707428 --- /dev/null +++ b/proto-cuda/packs/igneum-hourly/kernel.cl @@ -0,0 +1,176 @@ +// Generated by proto-metal/igneum-bench --export-pack for seed "igneum-hourly". Do not edit by hand. +// OpenCL C twin of the Metal kernel for the same seed (see proto-opencl/README.md, WAVEFRONT.md and program.metal). +// Built from source at runtime by proto-opencl/host.c, which passes these defines: +// IGNEUM_GROUP work-group size of igneum_hash, a multiple of 32 (default 32: one work-group = one 32-lane unit) +// IGNEUM_EXCHANGE 0 = local-memory exchange with a barrier (any device, any wave width; the default) +// 1 = sub_group_shuffle_xor (cl_khr_subgroup_shuffle), only with IGNEUM_GROUP 32 and a sub-group size of exactly 32 +// 2 = intel_sub_group_shuffle_xor (cl_intel_subgroups), same condition +// The verification unit is always 32 lanes. A 64-wide hardware wave (AMD GCN/CDNA, RDNA in wave64) runs two units; +// the exchange masks are 1, 2, 4, 8, 16, so every partner lane lies inside the lane's own aligned run of 32. +#ifndef IGNEUM_GROUP +#define IGNEUM_GROUP 32 +#endif +#ifndef IGNEUM_EXCHANGE +#define IGNEUM_EXCHANGE 0 +#endif +#ifdef __OPENCL_VERSION__ +#define IGNEUM_KERNEL_HASH __kernel __attribute__((reqd_work_group_size(IGNEUM_GROUP, 1, 1))) +#define IGNEUM_LOCAL_WORDS(name, n) __local uint name[n] +#if IGNEUM_EXCHANGE == 1 +#ifdef cl_khr_subgroups +#pragma OPENCL EXTENSION cl_khr_subgroups : enable +#endif +#ifdef cl_khr_subgroup_shuffle +#pragma OPENCL EXTENSION cl_khr_subgroup_shuffle : enable +#endif +#elif IGNEUM_EXCHANGE == 2 +#pragma OPENCL EXTENSION cl_intel_subgroups : enable +#endif +#else +// Not an OpenCL compiler: proto-opencl/emu compiles this file as C++ and supplies the built-ins and these two macros. +#include "emu_opencl.h" +#endif + +#if IGNEUM_EXCHANGE == 1 +#define IGNEUM_SHFL_XOR(dst, a, m) dst = sub_group_shuffle_xor((a), (uint)(m)) +#define IGNEUM_BCAST0(dst, a) dst = sub_group_broadcast((a), 0u) +#elif IGNEUM_EXCHANGE == 2 +#define IGNEUM_SHFL_XOR(dst, a, m) dst = intel_sub_group_shuffle_xor((a), (uint)(m)) +#define IGNEUM_BCAST0(dst, a) dst = sub_group_broadcast((a), 0u) +#else +// Local-memory exchange. Two buffers of IGNEUM_GROUP words alternate (xk counts exchanges), so one barrier per +// exchange is enough: a lane can only overwrite buffer b at exchange k+2 after passing barrier k+1, and every lane +// reaches barrier k+1 only after its read of buffer b at exchange k. The partner lid ^ m stays inside the lane's +// aligned run of 32 because m < 32. Control flow is uniform, so every work-item reaches every barrier. +#define IGNEUM_SHFL_XOR(dst, a, m) { xch[(xk & 1u) * IGNEUM_GROUP + lid] = (a); barrier(CLK_LOCAL_MEM_FENCE); dst = xch[(xk & 1u) * IGNEUM_GROUP + (lid ^ (uint)(m))]; xk += 1u; } +#define IGNEUM_BCAST0(dst, a) { xch[(xk & 1u) * IGNEUM_GROUP + lid] = (a); barrier(CLK_LOCAL_MEM_FENCE); dst = xch[(xk & 1u) * IGNEUM_GROUP + (lid & ~31u)]; xk += 1u; } +#endif + +static inline uint splitmix32(uint x) { + x ^= x >> 16; x *= 0x7feb352du; + x ^= x >> 15; x *= 0x846ca68bu; + x ^= x >> 16; + return x; +} +// n is a literal in 1..31 at every call site. OpenCL rotate() rotates left by n modulo 32. +static inline uint rotl_imm(uint x, uint n) { return rotate(x, n); } +// Right rotation by n modulo 32 as a left rotation by (32 - n) modulo 32; n == 0 gives x. +static inline uint rotr_var(uint x, uint n) { return rotate(x, (0u - n) & 31u); } +static inline uint ds_elem(uint i, uint d0, uint d1) { + uint x = i ^ d0; + x *= 0x9E3779B1u; x ^= x >> 15; + x += d1; + x *= 0x85EBCA77u; x ^= x >> 13; + x *= 0xC2B2AE3Du; x ^= x >> 16; + return x; +} + +// dataset[i] = ds_elem(i, d0, d1) for i < n. Same closed form as the Metal igneum_fill kernel. +__kernel void igneum_fill(__global uint* ds, uint n, uint d0, uint d1) { + uint i = (uint)get_global_id(0); + if (i < n) ds[i] = ds_elem(i, d0, d1); +} + +// One hash per work-item. IGNEUM_GROUP is a multiple of 32; lane = lid & 31 and every exchange stays inside the +// lane's own aligned run of 32 work-items, exactly like simd_shuffle_xor inside a 32-wide Metal SIMD group and +// __shfl_xor_sync inside a CUDA warp. Control flow is uniform (no branches at all). +IGNEUM_KERNEL_HASH void igneum_hash(__global const uint* ds, __global ulong* out, uint baseNonce, uint mask) { + uint gid = (uint)get_global_id(0); + uint lid = (uint)get_local_id(0); + uint nonce = baseNonce + gid; + uint r0, r1, r2, r3, r4, r5, r6, r7; +#if IGNEUM_EXCHANGE == 0 + IGNEUM_LOCAL_WORDS(xch, 2 * IGNEUM_GROUP); + uint xk = 0u; +#else + (void)lid; +#endif + { uint x = nonce ^ 0x6bdee811u; x += 0x9e3779b9u; x = splitmix32(x); r0 = x ^ 0x8f488bbeu; } // SEEDW[0], 0x9e3779b9u * 1u, SEEDW[1] + { uint x = nonce ^ 0x8f488bbeu; x += 0x3c6ef372u; x = splitmix32(x); r1 = x ^ 0xc5cdece7u; } // SEEDW[1], 0x9e3779b9u * 2u, SEEDW[2] + { uint x = nonce ^ 0xc5cdece7u; x += 0xdaa66d2bu; x = splitmix32(x); r2 = x ^ 0x210af22du; } // SEEDW[2], 0x9e3779b9u * 3u, SEEDW[3] + { uint x = nonce ^ 0x210af22du; x += 0x78dde6e4u; x = splitmix32(x); r3 = x ^ 0x2f687b65u; } // SEEDW[3], 0x9e3779b9u * 4u, SEEDW[4] + { uint x = nonce ^ 0x2f687b65u; x += 0x1715609du; x = splitmix32(x); r4 = x ^ 0x17471eeeu; } // SEEDW[4], 0x9e3779b9u * 5u, SEEDW[5] + { uint x = nonce ^ 0x17471eeeu; x += 0xb54cda56u; x = splitmix32(x); r5 = x ^ 0xee16e284u; } // SEEDW[5], 0x9e3779b9u * 6u, SEEDW[6] + { uint x = nonce ^ 0xee16e284u; x += 0x5384540fu; x = splitmix32(x); r6 = x ^ 0xfc9eb8f9u; } // SEEDW[6], 0x9e3779b9u * 7u, SEEDW[7] + { uint x = nonce ^ 0xfc9eb8f9u; x += 0xf1bbcdc8u; x = splitmix32(x); r7 = x ^ 0x6bdee811u; } // SEEDW[7], 0x9e3779b9u * 8u, SEEDW[0] + + for (uint it = 0u; it < 8u; ++it) { + uint sel = r0; + r6 = r6 ^ ds[r5 & mask]; // 0 load + r3 = r3 ^ r7; // 1 xor + r6 = mul_hi(r6, r2); // 2 mulhi + r1 = r1 + r0 + ((((sel >> 0u) & 1u) != 0u) ? 0x43f8f369u : 0x1eb46b1cu); // 3 add + r3 = r3 ^ ds[r0 & mask]; // 4 load + r5 = r5 | r7; // 5 or + r4 = r4 ^ ds[r6 & mask]; // 6 load + r4 = rotl_imm(r4, 21u); // 7 rotl + r6 = r6 ^ ds[r3 & mask]; // 8 load + r6 = r6 ^ ds[r1 & mask]; // 9 load + { uint t_; IGNEUM_SHFL_XOR(t_, r3, 8u); r0 = r0 ^ t_; } // 10 shfl + r2 = r2 ^ r3; // 11 xor + r2 = r2 + r7 + ((((sel >> 19u) & 1u) != 0u) ? 0xc26c7c2au : 0x3a1ce85eu); // 12 add + r4 = r4 ^ ds[r3 & mask]; // 13 load + r7 = r7 ^ ds[r1 & mask]; // 14 load + r2 = mul_hi(r2, r3); // 15 mulhi + r5 = r5 ^ ds[r2 & mask]; // 16 load + r5 = r5 ^ ds[r1 & mask]; // 17 load + r4 = r4 ^ ds[r7 & mask]; // 18 load + r2 = rotr_var(r2, r1); // 19 rotr + r7 = r7 ^ ds[r0 & mask]; // 20 load + r4 = r4 ^ ds[r6 & mask]; // 21 load + r7 = rotl_imm(r7, 25u); // 22 rotl + r3 = r3 + r5 + ((((sel >> 15u) & 1u) != 0u) ? 0x98ae0055u : 0x942b819bu); // 23 add + r3 = mul_hi(r3, r4); // 24 mulhi + r6 = r6 ^ r1; // 25 xor + r1 = rotl_imm(r1, 31u); // 26 rotl + { uint t_; IGNEUM_SHFL_XOR(t_, r6, 4u); r3 = r3 ^ t_; } // 27 shfl + r6 = r6 - r5; // 28 sub + r6 = rotr_var(r6, r3); // 29 rotr + { uint t_; IGNEUM_SHFL_XOR(t_, r4, 4u); r0 = r0 ^ t_; } // 30 shfl + r4 = rotl_imm(r4, 30u); // 31 rotl + r2 = r2 - r1; // 32 sub + r5 = r5 | r4; // 33 or + r7 = r6 * r3 + r7; // 34 mad + r5 = r5 * r0; // 35 mul + r5 = r5 - r3; // 36 sub + r2 = r2 + r7 + ((((sel >> 5u) & 1u) != 0u) ? 0x6b5970b5u : 0x473ecfd5u); // 37 add + r2 = r2 ^ ds[r7 & mask]; // 38 load + r2 = rotr_var(r2, r6); // 39 rotr + r0 = r0 ^ r5; // 40 xor + { uint t_; IGNEUM_SHFL_XOR(t_, r5, 8u); r4 = r4 ^ t_; } // 41 shfl + r1 = r1 * r6; // 42 mul + r0 = r4 * r1 + r0; // 43 mad + r1 = r1 + r4 + ((((sel >> 20u) & 1u) != 0u) ? 0x4a502c22u : 0x04e78f3bu); // 44 add + r6 = r2 * r7 + r6; // 45 mad + r1 = r1 ^ r0; // 46 xor + r5 = r5 ^ ds[r7 & mask]; // 47 load + r0 = r0 | r4; // 48 or + { uint t_; IGNEUM_SHFL_XOR(t_, r4, 16u); r5 = r5 ^ t_; } // 49 shfl + r7 = r7 + r0 + ((((sel >> 17u) & 1u) != 0u) ? 0x6fabf9ceu : 0x0a3df170u); // 50 add + r6 = r6 ^ r3; // 51 xor + r1 = r1 + r2 + ((((sel >> 21u) & 1u) != 0u) ? 0xad344ca0u : 0xc99bce6fu); // 52 add + r6 = r6 * r7; // 53 mul + r3 = mul_hi(r3, r4); // 54 mulhi + r7 = r7 * r1; // 55 mul + r7 = r7 ^ r1; // 56 xor + r2 = r2 * r7; // 57 mul + r2 = r2 ^ ds[r1 & mask]; // 58 load + r7 = r4 * r5 + r7; // 59 mad + r2 = r2 ^ ds[r7 & mask]; // 60 load + r0 = r2 * r3 + r0; // 61 mad + r1 = mul_hi(r1, r5); // 62 mulhi + r7 = r7 + r3 + ((((sel >> 24u) & 1u) != 0u) ? 0x05e4fc1du : 0xc23e27c9u); // 63 add + } + uint lo = r0 ^ rotl_imm(r1, 7u) ^ rotl_imm(r2, 14u) ^ rotl_imm(r3, 21u); + uint hi = r4 ^ rotl_imm(r5, 9u) ^ rotl_imm(r6, 18u) ^ rotl_imm(r7, 27u); + out[gid] = ((ulong)hi << 32) | (ulong)lo; +} + +#if IGNEUM_EXCHANGE != 0 +// Reports the sub-group size this device uses for a work-group of IGNEUM_GROUP items. host.c runs it only when the +// per-kernel query (clGetKernelSubGroupInfoKHR on igneum_hash) is unavailable; that query is preferred because a +// compiler may pick a different wave width per kernel (RDNA: wave32 or wave64). See WAVEFRONT.md. +IGNEUM_KERNEL_HASH void igneum_probe_subgroup(__global uint* out) { + if (get_local_id(0) == 0u) { out[0] = get_sub_group_size(); out[1] = get_num_sub_groups(); } +} +#endif diff --git a/proto-cuda/packs/igneum-hourly/program.h b/proto-cuda/packs/igneum-hourly/program.h index 4ac7453cd..0b7472f8c 100644 --- a/proto-cuda/packs/igneum-hourly/program.h +++ b/proto-cuda/packs/igneum-hourly/program.h @@ -1,8 +1,15 @@ // Generated by proto-metal/igneum-bench --export-pack for seed "igneum-hourly". Do not edit by hand. // Program metadata for host.cu plus the launch wrappers defined in kernel.cu. +// Also included by proto-opencl/host.c (C99), which defines IGNEUM_NO_CUDA first and reads only the macros. #pragma once -#include +#ifdef __cplusplus #include +#else +#include +#endif +#ifndef IGNEUM_NO_CUDA +#include +#endif #define IGNEUM_SEED_STRING "igneum-hourly" #define IGNEUM_DAY_STRING "2026-10-03" @@ -14,12 +21,16 @@ #define IGNEUM_ITERATIONS 8 #define IGNEUM_INSTR_COUNT 64 #define IGNEUM_LOADS_PER_HASH 128 +#define IGNEUM_WIDE_LOADS_PER_HASH 0 #define IGNEUM_OP_MIX "load=16 add=8 xor=7 mad=5 mul=5 mulhi=5 shfl=5 rotl=4 or=3 rotr=3 sub=3" +// 0 = closed-form dataset (ds_elem), 1 = memory-hard cache construction (MEMHARD.md, memhard.h) +#define IGNEUM_DATASET_MODE 0 #define IGNEUM_SEEDW_INIT { 0x6bdee811u, 0x8f488bbeu, 0xc5cdece7u, 0x210af22du, 0x2f687b65u, 0x17471eeeu, 0xee16e284u, 0xfc9eb8f9u } - +#ifndef IGNEUM_NO_CUDA // Defined in kernel.cu. Both launch on the default stream and return cudaGetLastError(). cudaError_t igneum_launch_fill(uint32_t* ds, uint32_t nWords, uint32_t d0, uint32_t d1); cudaError_t igneum_launch_hash(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask, uint32_t nonces, uint32_t blockWarps); cudaError_t igneum_hash_info(int* numRegs, int* blocksPerSM, uint32_t blockWarps); +#endif diff --git a/proto-cuda/packs/igneum-hourly/program.json b/proto-cuda/packs/igneum-hourly/program.json index 975de5d82..78e7066a0 100644 --- a/proto-cuda/packs/igneum-hourly/program.json +++ b/proto-cuda/packs/igneum-hourly/program.json @@ -1,5 +1,6 @@ { - "format": "igneum-program-pack-1", + "format": "igneum-program-pack-2", + "dataset_mode": "closed-form", "seed": "igneum-hourly", "seed_words": ["0x6bdee811", "0x8f488bbe", "0xc5cdece7", "0x210af22d", "0x2f687b65", "0x17471eee", "0xee16e284", "0xfc9eb8f9"], "seed_derivation": "FNV-1a 64 over UTF-8 of seed, basis ^ (salt * 0x9E3779B97F4A7C15) for salt 0..3, then h ^= h>>33; h *= 0xff51afd7ed558ccd; h ^= h>>33; words[2*salt] = low 32, words[2*salt+1] = high 32", @@ -24,7 +25,8 @@ "rotr": "dst = rotr(dst, src & 31)", "mad": "dst = src * src2 + dst", "shfl": "dst = dst ^ (src of lane (lane ^ mask)), mask in {1,2,4,8,16}, within the 32-lane warp", - "load": "dst = dst ^ dataset[src & dataset.mask]" + "load": "dst = dst ^ dataset[src & dataset.mask]", + "wload": "base = (src of lane 0 & dataset.mask) & ~31; dst = dst ^ dataset[base + lane] (warp-coalesced 128-byte load, lever b, only when --wide-frac > 0)" }, "dataset": { "log2_words": 28, @@ -34,6 +36,7 @@ "day_words_from": "day/2026-10-03", "d0": "0x3067619f", "d1": "0x3c269176", + "mode": "closed-form", "formula": "x = i ^ d0; x *= 0x9E3779B1; x ^= x>>15; x += d1; x *= 0x85EBCA77; x ^= x>>13; x *= 0xC2B2AE3D; x ^= x>>16 (all mod 2^32)" }, "instructions": [ diff --git a/proto-cuda/packs/igneum-hourly/vectors.h b/proto-cuda/packs/igneum-hourly/vectors.h index ed6c31f4d..d99665676 100644 --- a/proto-cuda/packs/igneum-hourly/vectors.h +++ b/proto-cuda/packs/igneum-hourly/vectors.h @@ -1,7 +1,11 @@ // Generated by proto-metal/igneum-bench --export-pack for seed "igneum-hourly". Do not edit by hand. -// Expected outputs: proto-metal CPU interpreter (cpuWarp) on Apple M5 Max; Metal GPU cross-check PASS 3/3 warps +// Expected outputs: proto-metal CPU interpreter (cpuWarp, closed-form dataset) on Apple M5 Max; Metal GPU cross-check PASS 3/3 warps #pragma once +#ifdef __cplusplus #include +#else +#include +#endif #define IGNEUM_VEC_WARPS 3 static const uint32_t IGNEUM_VEC_BASE[IGNEUM_VEC_WARPS] = { 0u, 4096u, 1000000u }; @@ -33,3 +37,11 @@ static const uint32_t IGNEUM_DS_HEAD[16] = { }; static const uint32_t IGNEUM_DS_LAST_INDEX = 268435455u; static const uint32_t IGNEUM_DS_LAST = 0xf78c84a4u; +// 64 sampled dataset words (index, value) computed on the Mac. +#define IGNEUM_DS_SAMPLES 64 +static const uint32_t IGNEUM_DS_SAMPLE_INDEX[IGNEUM_DS_SAMPLES] = { + 59471966u, 217795994u, 208353206u, 42483309u, 172547758u, 148076330u, 183853158u, 214389424u, 267488061u, 169781097u, 184093494u, 153880993u, 84977930u, 46426879u, 3093825u, 225364072u, 44593546u, 260713159u, 168250303u, 52384140u, 223401610u, 45554030u, 95410555u, 175039924u, 79171087u, 267580473u, 24168642u, 37981670u, 171551130u, 195559979u, 204611762u, 140997658u, 138925853u, 86637313u, 20736778u, 219665210u, 160430336u, 264654675u, 8013395u, 228945585u, 213884386u, 104419827u, 44185464u, 142737231u, 99284897u, 132475900u, 61861762u, 132056166u, 262388043u, 91878046u, 117353561u, 124768597u, 71352993u, 190698941u, 46055428u, 55281366u, 165145231u, 106810753u, 171985651u, 232085256u, 159510492u, 40072060u, 209107596u, 39023794u +}; +static const uint32_t IGNEUM_DS_SAMPLE_VALUE[IGNEUM_DS_SAMPLES] = { + 0x3ef2f18du, 0xed632a6bu, 0x2daad897u, 0x1201639fu, 0x8e79ca54u, 0x488d4f3du, 0x5206fa5cu, 0xf60c5f15u, 0x5af790f7u, 0xdb48b2cau, 0x2b55a0e4u, 0xc8c5100fu, 0xe4afb4a4u, 0x35a42586u, 0xff4cb6d1u, 0x1d6cac9fu, 0x889230c1u, 0x95316e2au, 0x61b099e9u, 0x97b89bb2u, 0x05171be5u, 0x7f0157abu, 0x13825368u, 0x562ab1bfu, 0x6e3b06b8u, 0xc9b51434u, 0xc994b628u, 0x26643aa8u, 0x56ae558cu, 0xdc8fdaaau, 0x2ed21c3fu, 0x3a29b91bu, 0xfd98ca1eu, 0x44c7dd8au, 0xaab155c8u, 0x0db02165u, 0x637c813au, 0x5648617eu, 0x3e163337u, 0x64bf0ad7u, 0x43381384u, 0x824dca58u, 0xab7bf62eu, 0x10880075u, 0x292498dau, 0x925fea4du, 0x606466a0u, 0x37923fc4u, 0x3544915bu, 0x742b205du, 0x3228e7adu, 0x1ce5ec92u, 0xd7857a5cu, 0x200a560bu, 0xc29076c7u, 0x886a1a06u, 0x1f415038u, 0x60fa6028u, 0x1fb1d0f9u, 0xee203b87u, 0xa12f7ea5u, 0x6ecd91dcu, 0x260e00b4u, 0xef4a4067u +}; diff --git a/proto-cuda/packs/igneum-hourly/vectors.json b/proto-cuda/packs/igneum-hourly/vectors.json index 9539f0b06..9c573d930 100644 --- a/proto-cuda/packs/igneum-hourly/vectors.json +++ b/proto-cuda/packs/igneum-hourly/vectors.json @@ -1,10 +1,11 @@ { "seed": "igneum-hourly", "day": "2026-10-03", + "dataset_mode": "closed-form", "dataset_log2_words": 28, "mask": "0x0fffffff", "lanes": 32, - "source": "proto-metal CPU interpreter (cpuWarp) on Apple M5 Max; Metal GPU cross-check PASS 3/3 warps", + "source": "proto-metal CPU interpreter (cpuWarp, closed-form dataset) on Apple M5 Max; Metal GPU cross-check PASS 3/3 warps", "warps": [ {"base_nonce": 0, "expected": [ "0x787a737506455bbe", "0xf93d248e364f7e1e", "0x8a7ce1701b5d0489", "0xd4f5b2177c414f3f", "0xe88fdf6fb22aee5c", "0xd68bc0baf30b41fc", "0x87f6a81eb0a6ac69", "0xb3cd8b6e6e57b347", @@ -27,5 +28,6 @@ ], "dataset_head": ["0x82174c0f", "0x577bdb9c", "0x111053bf", "0x2bb85514", "0xd83e190f", "0x8ccd9427", "0x3f69d5e4", "0xf3bbe1dc", "0xb3aa7e90", "0xc33f5f73", "0xb83a2b10", "0xc4d7c8ff", "0xefa6d1a8", "0x7029b116", "0x5e48fab0", "0x0ef66e40"], "dataset_last_index": 268435455, - "dataset_last": "0xf78c84a4" + "dataset_last": "0xf78c84a4", + "dataset_samples": [{"index": 59471966, "value": "0x3ef2f18d"}, {"index": 217795994, "value": "0xed632a6b"}, {"index": 208353206, "value": "0x2daad897"}, {"index": 42483309, "value": "0x1201639f"}, {"index": 172547758, "value": "0x8e79ca54"}, {"index": 148076330, "value": "0x488d4f3d"}, {"index": 183853158, "value": "0x5206fa5c"}, {"index": 214389424, "value": "0xf60c5f15"}, {"index": 267488061, "value": "0x5af790f7"}, {"index": 169781097, "value": "0xdb48b2ca"}, {"index": 184093494, "value": "0x2b55a0e4"}, {"index": 153880993, "value": "0xc8c5100f"}, {"index": 84977930, "value": "0xe4afb4a4"}, {"index": 46426879, "value": "0x35a42586"}, {"index": 3093825, "value": "0xff4cb6d1"}, {"index": 225364072, "value": "0x1d6cac9f"}, {"index": 44593546, "value": "0x889230c1"}, {"index": 260713159, "value": "0x95316e2a"}, {"index": 168250303, "value": "0x61b099e9"}, {"index": 52384140, "value": "0x97b89bb2"}, {"index": 223401610, "value": "0x05171be5"}, {"index": 45554030, "value": "0x7f0157ab"}, {"index": 95410555, "value": "0x13825368"}, {"index": 175039924, "value": "0x562ab1bf"}, {"index": 79171087, "value": "0x6e3b06b8"}, {"index": 267580473, "value": "0xc9b51434"}, {"index": 24168642, "value": "0xc994b628"}, {"index": 37981670, "value": "0x26643aa8"}, {"index": 171551130, "value": "0x56ae558c"}, {"index": 195559979, "value": "0xdc8fdaaa"}, {"index": 204611762, "value": "0x2ed21c3f"}, {"index": 140997658, "value": "0x3a29b91b"}, {"index": 138925853, "value": "0xfd98ca1e"}, {"index": 86637313, "value": "0x44c7dd8a"}, {"index": 20736778, "value": "0xaab155c8"}, {"index": 219665210, "value": "0x0db02165"}, {"index": 160430336, "value": "0x637c813a"}, {"index": 264654675, "value": "0x5648617e"}, {"index": 8013395, "value": "0x3e163337"}, {"index": 228945585, "value": "0x64bf0ad7"}, {"index": 213884386, "value": "0x43381384"}, {"index": 104419827, "value": "0x824dca58"}, {"index": 44185464, "value": "0xab7bf62e"}, {"index": 142737231, "value": "0x10880075"}, {"index": 99284897, "value": "0x292498da"}, {"index": 132475900, "value": "0x925fea4d"}, {"index": 61861762, "value": "0x606466a0"}, {"index": 132056166, "value": "0x37923fc4"}, {"index": 262388043, "value": "0x3544915b"}, {"index": 91878046, "value": "0x742b205d"}, {"index": 117353561, "value": "0x3228e7ad"}, {"index": 124768597, "value": "0x1ce5ec92"}, {"index": 71352993, "value": "0xd7857a5c"}, {"index": 190698941, "value": "0x200a560b"}, {"index": 46055428, "value": "0xc29076c7"}, {"index": 55281366, "value": "0x886a1a06"}, {"index": 165145231, "value": "0x1f415038"}, {"index": 106810753, "value": "0x60fa6028"}, {"index": 171985651, "value": "0x1fb1d0f9"}, {"index": 232085256, "value": "0xee203b87"}, {"index": 159510492, "value": "0xa12f7ea5"}, {"index": 40072060, "value": "0x6ecd91dc"}, {"index": 209107596, "value": "0x260e00b4"}, {"index": 39023794, "value": "0xef4a4067"}] } diff --git a/proto-metal/main.swift b/proto-metal/main.swift index e8d5bd8db..6c7fbbcca 100644 --- a/proto-metal/main.swift +++ b/proto-metal/main.swift @@ -69,7 +69,7 @@ func parseArgs() -> Options { [--closed-form] original closed-form dataset (default: memory-hard cache construction, see MEMHARD.md) [--load-weight W] generator lever (a): percent weight of the load op (default 25) [--wide-frac P] generator lever (b): percent of loads emitted as warp-coalesced 128-byte loads (default 0) - [--export-pack ] write the CUDA program pack for --seed, then exit + [--export-pack ] write the CUDA + OpenCL program pack for --seed, then exit hardening tests (run instead of the bench; several may be combined; exit 0 only if all pass): [--fuzz N [--fuzz-seed ]] N random programs, GPU vs CPU, 4 random warps each, dataset size drawn from 64 MiB, 256 MiB, 1 GiB @@ -420,16 +420,21 @@ func generateProgram(seedString: String) -> Program { func hex(_ v: UInt32) -> String { String(format: "0x%08xu", v) } -// The memory-hard core as source text, in Metal (cuda false) or CUDA C++ (cuda true). Mixer parameters are -// literals so the GPU kernels, the CUDA pack and its host reference share one text. Names are prefixed mh_. +// The memory-hard core as source text, in Metal, CUDA C++ or OpenCL C. Mixer parameters are +// literals so the GPU kernels, the CUDA pack, the OpenCL pack and the host reference share one text. Names are prefixed mh_. // In the CUDA dialect every function is IGNEUM_HD (host and device) so host.cu can derive items too. -func emitMemhardCore(_ mp: MixParams, cuda: Bool) -> String { - let U = cuda ? "uint32_t" : "uint" - let fn = cuda ? "IGNEUM_HD" : "inline" - let cptr = cuda ? "const uint32_t*" : "device const uint*" - let wptr = cuda ? "uint32_t*" : "device uint*" - let lptr = cuda ? "uint32_t*" : "thread uint*" - let lcptr = cuda ? "const uint32_t*" : "const thread uint*" +// In the OpenCL dialect the cache pointers carry the __global address space (OpenCL C 1.2 has no generic space). +enum CoreDialect { case metal, cuda, opencl } + +func emitMemhardCore(_ mp: MixParams, cuda: Bool) -> String { emitMemhardCore(mp, dialect: cuda ? .cuda : .metal) } + +func emitMemhardCore(_ mp: MixParams, dialect: CoreDialect) -> String { + let U: String, fn: String, cptr: String, wptr: String, lptr: String, lcptr: String + switch dialect { + case .metal: (U, fn, cptr, wptr, lptr, lcptr) = ("uint", "inline", "device const uint*", "device uint*", "thread uint*", "const thread uint*") + case .cuda: (U, fn, cptr, wptr, lptr, lcptr) = ("uint32_t", "IGNEUM_HD", "const uint32_t*", "uint32_t*", "uint32_t*", "const uint32_t*") + case .opencl: (U, fn, cptr, wptr, lptr, lcptr) = ("uint", "static inline", "__global const uint*", "__global uint*", "uint*", "const uint*") + } let K = mp.keyWords, R = mp.rotWords, M = mp.mulWords, C = mp.rcWords var s = """ // Memory-hard dataset core (MEMHARD.md). Cache: 2^\(cacheLog2Words) words in 2^\(Int(log2(Double(cacheSegments)))) segments of \(cacheLinesPerSegment) chained ChaCha\(chachaRounds) lines. @@ -722,7 +727,7 @@ func cpuWarpTraced(_ p: Program, baseNonce: UInt32, ds: DatasetSource, // MARK: - Program pack export (CUDA twin of the Metal kernel) // -// Writes, for one seed: program.json, vectors.json, kernel.cu, program.h, vectors.h, program.metal. +// Writes, for one seed: program.json, vectors.json, kernel.cu, kernel.cl, program.h, vectors.h, program.metal. // The CUDA kernel is emitted from the same Instr list as the MSL above, line for line. // Differences by design: the dataset mask is a kernel argument (so the host can sweep dataset // sizes with one ahead-of-time compile), seeds are inlined as literals, and the host launch @@ -910,15 +915,192 @@ func generateCUDA(_ p: Program, memhard: MixParams?) -> String { return s } +// kernel.cl: the same program as OpenCL C 1.2, built from source at runtime by proto-opencl/host.c. +// The 32-lane exchange is selected at compile time (IGNEUM_EXCHANGE): sub-group shuffles where the device has them +// and its sub-group size is exactly 32, or a local-memory exchange with a barrier on every other device +// (proto-opencl/WAVEFRONT.md). The same text is compiled as C++ by proto-opencl/emu with a 32- or 64-wide sub-group. +func generateOpenCL(_ p: Program, memhard: MixParams?) -> String { + var s = """ + // Generated by proto-metal/igneum-bench --export-pack for seed "\(p.seedString)". Do not edit by hand. + // OpenCL C twin of the Metal kernel for the same seed (see proto-opencl/README.md, WAVEFRONT.md and program.metal). + // Built from source at runtime by proto-opencl/host.c, which passes these defines: + // IGNEUM_GROUP work-group size of igneum_hash, a multiple of 32 (default 32: one work-group = one 32-lane unit) + // IGNEUM_EXCHANGE 0 = local-memory exchange with a barrier (any device, any wave width; the default) + // 1 = sub_group_shuffle_xor (cl_khr_subgroup_shuffle), only with IGNEUM_GROUP 32 and a sub-group size of exactly 32 + // 2 = intel_sub_group_shuffle_xor (cl_intel_subgroups), same condition + // The verification unit is always 32 lanes. A 64-wide hardware wave (AMD GCN/CDNA, RDNA in wave64) runs two units; + // the exchange masks are 1, 2, 4, 8, 16, so every partner lane lies inside the lane's own aligned run of 32. + #ifndef IGNEUM_GROUP + #define IGNEUM_GROUP 32 + #endif + #ifndef IGNEUM_EXCHANGE + #define IGNEUM_EXCHANGE 0 + #endif + #ifdef __OPENCL_VERSION__ + #define IGNEUM_KERNEL_HASH __kernel __attribute__((reqd_work_group_size(IGNEUM_GROUP, 1, 1))) + #define IGNEUM_LOCAL_WORDS(name, n) __local uint name[n] + #if IGNEUM_EXCHANGE == 1 + #ifdef cl_khr_subgroups + #pragma OPENCL EXTENSION cl_khr_subgroups : enable + #endif + #ifdef cl_khr_subgroup_shuffle + #pragma OPENCL EXTENSION cl_khr_subgroup_shuffle : enable + #endif + #elif IGNEUM_EXCHANGE == 2 + #pragma OPENCL EXTENSION cl_intel_subgroups : enable + #endif + #else + // Not an OpenCL compiler: proto-opencl/emu compiles this file as C++ and supplies the built-ins and these two macros. + #include "emu_opencl.h" + #endif + + #if IGNEUM_EXCHANGE == 1 + #define IGNEUM_SHFL_XOR(dst, a, m) dst = sub_group_shuffle_xor((a), (uint)(m)) + #define IGNEUM_BCAST0(dst, a) dst = sub_group_broadcast((a), 0u) + #elif IGNEUM_EXCHANGE == 2 + #define IGNEUM_SHFL_XOR(dst, a, m) dst = intel_sub_group_shuffle_xor((a), (uint)(m)) + #define IGNEUM_BCAST0(dst, a) dst = sub_group_broadcast((a), 0u) + #else + // Local-memory exchange. Two buffers of IGNEUM_GROUP words alternate (xk counts exchanges), so one barrier per + // exchange is enough: a lane can only overwrite buffer b at exchange k+2 after passing barrier k+1, and every lane + // reaches barrier k+1 only after its read of buffer b at exchange k. The partner lid ^ m stays inside the lane's + // aligned run of 32 because m < 32. Control flow is uniform, so every work-item reaches every barrier. + #define IGNEUM_SHFL_XOR(dst, a, m) { xch[(xk & 1u) * IGNEUM_GROUP + lid] = (a); barrier(CLK_LOCAL_MEM_FENCE); dst = xch[(xk & 1u) * IGNEUM_GROUP + (lid ^ (uint)(m))]; xk += 1u; } + #define IGNEUM_BCAST0(dst, a) { xch[(xk & 1u) * IGNEUM_GROUP + lid] = (a); barrier(CLK_LOCAL_MEM_FENCE); dst = xch[(xk & 1u) * IGNEUM_GROUP + (lid & ~31u)]; xk += 1u; } + #endif + + static inline uint splitmix32(uint x) { + x ^= x >> 16; x *= 0x7feb352du; + x ^= x >> 15; x *= 0x846ca68bu; + x ^= x >> 16; + return x; + } + // n is a literal in 1..31 at every call site. OpenCL rotate() rotates left by n modulo 32. + static inline uint rotl_imm(uint x, uint n) { return rotate(x, n); } + // Right rotation by n modulo 32 as a left rotation by (32 - n) modulo 32; n == 0 gives x. + static inline uint rotr_var(uint x, uint n) { return rotate(x, (0u - n) & 31u); } + static inline uint ds_elem(uint i, uint d0, uint d1) { + uint x = i ^ d0; + x *= 0x9E3779B1u; x ^= x >> 15; + x += d1; + x *= 0x85EBCA77u; x ^= x >> 13; + x *= 0xC2B2AE3Du; x ^= x >> 16; + return x; + } + + + """ + if let mp = memhard { + s += emitMemhardCore(mp, dialect: .opencl) + "\n" + s += """ + // Memory-hard dataset (MEMHARD.md). One work-item per cache segment; one work-item per 64-byte dataset item. + // The same constants as memhard.h in this pack (one emitter, three dialects). + __kernel void igneum_cache_fill(__global uint* cache, uint nSegments) { + uint seg = (uint)get_global_id(0); + if (seg < nSegments) mh_cache_segment(cache, seg); + } + __kernel void igneum_build(__global uint* ds, __global const uint* cache, uint nItems) { + uint t = (uint)get_global_id(0); + if (t < nItems) { + uint s[16]; + mh_item(cache, t, s); + __global uint* d = ds + ((ulong)t * 16u); + for (uint i = 0u; i < 16u; ++i) d[i] = s[i]; + } + } + + + """ + } else { + s += """ + // dataset[i] = ds_elem(i, d0, d1) for i < n. Same closed form as the Metal igneum_fill kernel. + __kernel void igneum_fill(__global uint* ds, uint n, uint d0, uint d1) { + uint i = (uint)get_global_id(0); + if (i < n) ds[i] = ds_elem(i, d0, d1); + } + + + """ + } + s += """ + // One hash per work-item. IGNEUM_GROUP is a multiple of 32; lane = lid & 31 and every exchange stays inside the + // lane's own aligned run of 32 work-items, exactly like simd_shuffle_xor inside a 32-wide Metal SIMD group and + // __shfl_xor_sync inside a CUDA warp. Control flow is uniform (no branches at all). + IGNEUM_KERNEL_HASH void igneum_hash(__global const uint* ds, __global ulong* out, uint baseNonce, uint mask) { + uint gid = (uint)get_global_id(0); + uint lid = (uint)get_local_id(0); + uint nonce = baseNonce + gid; + uint r0, r1, r2, r3, r4, r5, r6, r7; + #if IGNEUM_EXCHANGE == 0 + IGNEUM_LOCAL_WORDS(xch, 2 * IGNEUM_GROUP); + uint xk = 0u; + #else + (void)lid; + #endif + + """ + if p.hasWide { s += " uint lane = lid & 31u;\n uint wmask = mask & ~31u;\n" } + for i in 0..<8 { + let addc = 0x9e3779b9 &* UInt32(i + 1) + s += " { uint x = nonce ^ \(hex(p.seed[i])); x += \(hex(addc)); x = splitmix32(x); r\(i) = x ^ \(hex(p.seed[(i + 1) & 7])); } // SEEDW[\(i)], 0x9e3779b9u * \(i + 1)u, SEEDW[\((i + 1) & 7)]\n" + } + s += "\n for (uint it = 0u; it < \(Program.iterations)u; ++it) {\n uint sel = r0;\n" + for (k, ins) in p.instrs.enumerated() { + let d = "r\(ins.dst)", a = "r\(ins.a)", b = "r\(ins.b)" + var line: String + switch ins.op { + // Metal select(A, B, c) returns c ? B : A, so the true branch is imm2 here as well. + case .add: line = "\(d) = \(d) + \(a) + ((((sel >> \(ins.bit)u) & 1u) != 0u) ? \(hex(ins.imm2)) : \(hex(ins.imm)));" + case .sub: line = "\(d) = \(d) - \(a);" + case .mul: line = "\(d) = \(d) * \(a);" + case .mulhi: line = "\(d) = mul_hi(\(d), \(a));" + case .xor: line = "\(d) = \(d) ^ \(a);" + case .or: line = "\(d) = \(d) | \(a);" + case .rotl: line = "\(d) = rotl_imm(\(d), \(ins.rot)u);" + case .rotr: line = "\(d) = rotr_var(\(d), \(a));" + case .mad: line = "\(d) = \(a) * \(b) + \(d);" + case .shfl: line = "{ uint t_; IGNEUM_SHFL_XOR(t_, \(a), \(ins.mask)u); \(d) = \(d) ^ t_; }" + case .load: line = "\(d) = \(d) ^ ds[\(a) & mask];" + case .wload: line = "{ uint t_; IGNEUM_BCAST0(t_, \(a)); \(d) = \(d) ^ ds[(t_ & wmask) + lane]; }" + } + s += " \(line) // \(k) \(ins.op.rawValue)\n" + } + s += """ + } + uint lo = r0 ^ rotl_imm(r1, 7u) ^ rotl_imm(r2, 14u) ^ rotl_imm(r3, 21u); + uint hi = r4 ^ rotl_imm(r5, 9u) ^ rotl_imm(r6, 18u) ^ rotl_imm(r7, 27u); + out[gid] = ((ulong)hi << 32) | (ulong)lo; + } + + #if IGNEUM_EXCHANGE != 0 + // Reports the sub-group size this device uses for a work-group of IGNEUM_GROUP items. host.c runs it only when the + // per-kernel query (clGetKernelSubGroupInfoKHR on igneum_hash) is unavailable; that query is preferred because a + // compiler may pick a different wave width per kernel (RDNA: wave32 or wave64). See WAVEFRONT.md. + IGNEUM_KERNEL_HASH void igneum_probe_subgroup(__global uint* out) { + if (get_local_id(0) == 0u) { out[0] = get_sub_group_size(); out[1] = get_num_sub_groups(); } + } + #endif + + """ + return s +} + func generateProgramHeader(_ p: Program, dayString: String, day: (UInt32, UInt32), datasetLog2: Int, memhard: MixParams?) -> String { let mask = UInt32((1 << datasetLog2) - 1) let mix = p.histogram.map { "\($0.0)=\($0.1)" }.joined(separator: " ") var s = """ // Generated by proto-metal/igneum-bench --export-pack for seed "\(p.seedString)". Do not edit by hand. // Program metadata for host.cu plus the launch wrappers defined in kernel.cu. + // Also included by proto-opencl/host.c (C99), which defines IGNEUM_NO_CUDA first and reads only the macros. #pragma once - #include + #ifdef __cplusplus #include + #else + #include + #endif + #ifndef IGNEUM_NO_CUDA + #include + #endif #define IGNEUM_SEED_STRING \(jstr(p.seedString)) #define IGNEUM_DAY_STRING \(jstr(dayString)) @@ -949,6 +1131,7 @@ func generateProgramHeader(_ p: Program, dayString: String, day: (UInt32, UInt32 #define IGNEUM_MIX_MUL_INIT { \(mp.mulWords.map(hex).joined(separator: ", ")) } #define IGNEUM_MIX_RC_INIT { \(mp.rcWords.map(hex).joined(separator: ", ")) } + #ifndef IGNEUM_NO_CUDA // Defined in kernel.cu. All launch on the default stream and return cudaGetLastError(). cudaError_t igneum_launch_cache_fill(uint32_t* cache, uint32_t nSegments); cudaError_t igneum_launch_build(uint32_t* ds, const uint32_t* cache, uint32_t nItems); @@ -956,6 +1139,7 @@ func generateProgramHeader(_ p: Program, dayString: String, day: (UInt32, UInt32 """ } else { s += """ + #ifndef IGNEUM_NO_CUDA // Defined in kernel.cu. Both launch on the default stream and return cudaGetLastError(). cudaError_t igneum_launch_fill(uint32_t* ds, uint32_t nWords, uint32_t d0, uint32_t d1); @@ -965,6 +1149,7 @@ func generateProgramHeader(_ p: Program, dayString: String, day: (UInt32, UInt32 cudaError_t igneum_launch_hash(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask, uint32_t nonces, uint32_t blockWarps); cudaError_t igneum_hash_info(int* numRegs, int* blocksPerSM, uint32_t blockWarps); + #endif """ return s @@ -975,15 +1160,22 @@ func generateMemhardHeader(_ p: Program, _ mp: MixParams) -> String { return """ // Generated by proto-metal/igneum-bench --export-pack for seed "\(p.seedString)". Do not edit by hand. // Memory-hard dataset core, the same text that the Mac's Metal kernels and CPU verifier were checked against. - // Included by kernel.cu (device) and host.cu (host reference). See proto-metal/MEMHARD.md for the construction. + // Included by kernel.cu (device), host.cu (host reference) and proto-opencl/host.c (C99 host reference). + // See proto-metal/MEMHARD.md for the construction. kernel.cl carries the same text in OpenCL C. #pragma once + #ifdef __cplusplus #include + #else + #include + #endif #if defined(__CUDACC__) #define IGNEUM_HD __host__ __device__ __forceinline__ + #elif defined(_MSC_VER) && !defined(__cplusplus) + #define IGNEUM_HD static __inline #else #define IGNEUM_HD static inline #endif - \(emitMemhardCore(mp, cuda: true)) + \(emitMemhardCore(mp, dialect: .cuda)) """ } @@ -1002,7 +1194,11 @@ func generateVectorsHeader(_ p: Program, bases: [UInt32], outs: [[UInt64]], v: P // Generated by proto-metal/igneum-bench --export-pack for seed "\(p.seedString)". Do not edit by hand. // Expected outputs: \(source) #pragma once + #ifdef __cplusplus #include + #else + #include + #endif #define IGNEUM_VEC_WARPS \(bases.count) static const uint32_t IGNEUM_VEC_BASE[IGNEUM_VEC_WARPS] = { \(bases.map { "\($0)u" }.joined(separator: ", ")) }; @@ -1204,6 +1400,7 @@ func exportPack(_ opts: Options) -> Never { ("program.json", generateProgramJSON(program, dayString: opts.day, day: day, datasetLog2: opts.datasetLog2, memhard: ctx.mp)), ("vectors.json", generateVectorsJSON(program, dayString: opts.day, datasetLog2: opts.datasetLog2, bases: packVectorBases, outs: outs, v: v, mask: mask, source: source, memhard: !ctx.closed)), ("kernel.cu", generateCUDA(program, memhard: ctx.mp)), + ("kernel.cl", generateOpenCL(program, memhard: ctx.mp)), ("program.h", generateProgramHeader(program, dayString: opts.day, day: day, datasetLog2: opts.datasetLog2, memhard: ctx.mp)), ("vectors.h", generateVectorsHeader(program, bases: packVectorBases, outs: outs, v: v, mask: mask, source: source, memhard: !ctx.closed)), ("program.metal", generateMSL(program, datasetLog2: opts.datasetLog2)), diff --git a/proto-opencl/.gitignore b/proto-opencl/.gitignore new file mode 100644 index 000000000..c0d505789 --- /dev/null +++ b/proto-opencl/.gitignore @@ -0,0 +1,6 @@ +igneum-bench-cl-* +*.exe +*.obj +*.o +emu/build-*/ +.DS_Store diff --git a/proto-opencl/README.md b/proto-opencl/README.md new file mode 100644 index 000000000..7e53418c1 --- /dev/null +++ b/proto-opencl/README.md @@ -0,0 +1,213 @@ +# igneum-bench-cl (proto-opencl) + +The portable third path for Igneum's random-program proof-of-work kernels, after Apple Metal (`proto-metal`) and +NVIDIA CUDA (`proto-cuda`). It runs the same program packs through OpenCL 1.2, which is what an AMD card exposes on +both Windows (Adrenalin driver) and Linux (ROCm, or Mesa), and checks every output bit for bit against the Mac. + +This is a test harness, not a miner. No pool, no network, no wallet, no mining protocol. It fills a dataset, checks the +device against known answers, and times the kernel. Nothing here earns anything. + +Status on 3 October 2026: no AMD device has run this yet. Everything that could be proven without one has been +(`WAVEFRONT.md`, "What was proven"): Apple's deprecated OpenCL 1.2 runtime on the M5 Max, pocl 7.2 on the Mac's CPU, +and a CPU emulator with 32- and 64-wide sub-groups all give the Mac's 96/96 vectors and cache FNV for the memory-hard +pack. The AMD run itself is the next step and the commands for it are below. + +## Layout + +``` +proto-opencl/ + host.c C99 host: device list, runtime kernel build, cache + dataset fill, self-tests, vectors, bench, sweep + build.sh macOS (-framework OpenCL, or the Khronos ICD loader) and Linux (-lOpenCL) + build.bat Windows (MSVC cl.exe + OpenCL.lib) + WAVEFRONT.md wave32 vs wave64 on AMD, and why the kernel cannot tell the difference + emu/ CPU emulator: compiles kernel.cl as C++ and runs it with a 32- or 64-wide sub-group +../proto-cuda/packs//kernel.cl the OpenCL C kernel of each pack, written by proto-metal/igneum-bench --export-pack +``` + +The packs stay in `proto-cuda/packs/` because the three implementations share one `program.h`, `vectors.h` and +`memhard.h` per seed; `kernel.cl` sits next to `kernel.cu` and `program.metal`. Three packs are checked in: +`igneum-genesis-mh` (memory-hard dataset, the current construction), `igneum-genesis` and `igneum-hourly` +(closed-form dataset, kept for comparison). + +## What the OpenCL path does differently + +| Item | CUDA (`host.cu`) | OpenCL (`host.c`) | +|---|---|---| +| Kernel compile | ahead of time by nvcc | at runtime by the driver from `kernel.cl`, with `-D IGNEUM_GROUP=32W -D IGNEUM_EXCHANGE=M` | +| 32-lane exchange | `__shfl_xor_sync` | `sub_group_shuffle_xor` only when the device's sub-group size is exactly 32 for a 32-item work-group; otherwise `__local` memory and a barrier (`WAVEFRONT.md`) | +| Timing | cudaEvent | event profiling (`CL_PROFILING_COMMAND_START/END`); wall time on Apple, whose OpenCL timestamps are unusable | +| Device choice | `--device D` | `--list`, then `--device D` by flat index over all platforms; default is the first GPU | +| Language | C++17 | C99, so MSVC builds it without nvcc or any C++ toolchain | + +Everything else (dataset fill or build, cache check against the Mac's FNV, dataset self-test, 3 vector warps standalone +and inside the warm-up batch, 5 timed batches of 2^24, the sweep, the summary table, exit codes) mirrors `host.cu` +line for line. Rates have the same definition: Mhash/s, and GB/s useful = loads per hash x 4 bytes x hashes/s. + +## Build + +Linux (AMD ROCm, AMD Adrenalin for Linux, or Mesa rusticl; any of them registers with the system ICD loader): + +``` +sudo apt install ocl-icd-opencl-dev opencl-headers clinfo # Debian or Ubuntu; Fedora: ocl-icd-devel opencl-headers clinfo +clinfo | head -40 # the AMD device must appear here before anything else matters +cd proto-opencl +./build.sh igneum-genesis-mh # -> ./igneum-bench-cl-igneum-genesis-mh +./build.sh igneum-genesis # closed-form packs +./build.sh igneum-hourly +``` + +The one command behind the script, for a pack `P`: + +``` +cc -std=c99 -O2 -Wall -Wextra -I ../proto-cuda/packs/P -DIGNEUM_KERNEL_PATH='"../proto-cuda/packs/P/kernel.cl"' -o igneum-bench-cl-P host.c -lOpenCL -ldl +``` + +Windows (AMD Adrenalin driver installed; it provides `OpenCL.dll` and the AMD ICD, nothing else is needed to run): + +1. Install Visual Studio 2022 or 2026 Build Tools with "Desktop development with C++". No CUDA, no nvcc. +2. Get the OpenCL headers and the import library `OpenCL.lib`, any one of these: + `vcpkg install opencl:x64-windows` then `set OPENCL_SDK=C:\vcpkg\installed\x64-windows`; + or the Khronos OpenCL-SDK release zip (GitHub KhronosGroup/OpenCL-SDK) unpacked anywhere, `set OPENCL_SDK=`; + or an installed CUDA Toolkit, which ships both under `%CUDA_PATH%` (the script finds it on its own). +3. From an "x64 Native Tools Command Prompt": + +``` +cd proto-opencl +build.bat igneum-genesis-mh +build.bat igneum-genesis +build.bat igneum-hourly +``` + +The one command behind the script: + +``` +cl /nologo /O2 /W3 /std:c11 /I "%OPENCL_SDK%\include" /I "..\proto-cuda\packs\P" /DIGNEUM_KERNEL_PATH="\"../proto-cuda/packs/P/kernel.cl\"" /Fe:igneum-bench-cl-P.exe host.c /link "%OPENCL_SDK%\lib\OpenCL.lib" +``` + +macOS (correctness check only; Apple deprecated OpenCL in macOS 10.14 and the runtime is 1.2): + +``` +./build.sh igneum-genesis-mh # Apple OpenCL.framework +./build.sh igneum-genesis-mh khr # Khronos ICD loader from Homebrew, for pocl: brew install opencl-headers opencl-icd-loader pocl +OCL_ICD_VENDORS=/opt/homebrew/etc/OpenCL/vendors SDKROOT=$(xcrun --show-sdk-path) ./igneum-bench-cl-igneum-genesis-mh-khr --list +``` + +(`SDKROOT` is needed because pocl links its kernels with the system linker and otherwise cannot find `-lSystem`.) + +## Run on the AMD rig + +Run from `proto-opencl/` so the default kernel path resolves, or pass `--kernel`. Exactly this, in this order, and paste +the whole stdout back: + +Windows: + +``` +cd proto-opencl +igneum-bench-cl-igneum-genesis-mh.exe --list +igneum-bench-cl-igneum-genesis-mh.exe > amd-mh-auto.txt +igneum-bench-cl-igneum-genesis-mh.exe --exchange local > amd-mh-local.txt +igneum-bench-cl-igneum-genesis-mh.exe --group-warps 2 > amd-mh-gw2.txt +igneum-bench-cl-igneum-genesis-mh.exe --sweep > amd-mh-sweep.txt +igneum-bench-cl-igneum-genesis.exe > amd-genesis.txt +igneum-bench-cl-igneum-hourly.exe > amd-hourly.txt +``` + +Linux: + +``` +cd proto-opencl +./igneum-bench-cl-igneum-genesis-mh --list +./igneum-bench-cl-igneum-genesis-mh | tee amd-mh-auto.txt +./igneum-bench-cl-igneum-genesis-mh --exchange local | tee amd-mh-local.txt +./igneum-bench-cl-igneum-genesis-mh --group-warps 2 | tee amd-mh-gw2.txt +./igneum-bench-cl-igneum-genesis-mh --sweep | tee amd-mh-sweep.txt +./igneum-bench-cl-igneum-genesis | tee amd-genesis.txt +./igneum-bench-cl-igneum-hourly | tee amd-hourly.txt +``` + +If `--list` shows more than one device (an iGPU, a CPU runtime, an NVIDIA card in the same box), add `--device D` +with the AMD card's index to every line. If the auto run says `exchange: sub_group_shuffle_xor`, also run +`--exchange subgroup` once so the log has an explicit sub-group run; if it says local memory, the `--exchange local` +run is a repeat and that is fine. + +Flags: `--list`, `--device D`, `--dataset-mib N` (power of two, default 1024), `--sweep` (4, 64, 256, 512, 1024 MiB), +`--batch-log2 B` (default 24), `--batches N` (default 5), `--group-warps W` (32-lane units per work-group, 1..8, +default 1), `--exchange auto|local|subgroup`, `--kernel path`, `--build-opts "..."` (appended to clBuildProgram), +`--time event|wall`. + +## What PASS looks like + +In order, the run prints: + +1. Every platform and device, with vendor, driver, OpenCL C version, compute units, memory, local memory, the + sub-group extension it lists, and (AMD) the wavefront width or (NVIDIA) the warp size. The chosen device is starred. +2. `build options:` and `exchange:`. The exchange line is the one that matters on AMD; it says which path was taken and + why, for example `local-memory exchange with barrier (the sub-group size for a 32-item work-group is not 32; queried + sub-group size 64)` on a wave64 card, or `sub_group_shuffle_xor (cl_khr_subgroup_shuffle), sub-group size 32 for a + 32-item work-group` on a wave32 one. Both are correct by construction (`WAVEFRONT.md`). +3. `kernel:` (work-group limit and local memory of `igneum_hash`), `program:`, `seed words:`, `day`. +4. Memory-hard packs: cache fill time on the device (twice) and on one host thread, then `cache check: PASS (device == + host all 67108864 words PASS, host FNV-1a 64 48c4f5bf24166b2e vs Mac 48c4f5bf24166b2e PASS, head 16 vs Mac PASS, + last line vs Mac PASS)`. The FNV is the same on every machine that has ever run this construction for day + 2026-10-03. +5. `dataset build` (or `dataset fill`) times, then `dataset self-test: PASS (head 16 vs Mac PASS, element [MASK] vs Mac + PASS, 64 random points vs host derivation PASS, 64 Mac samples PASS)`. +6. Six `verify warp ... PASS` lines: three standalone (bases 0, 4096, 1000000) and the same three read out of the + warm-up batch. +7. `batch fingerprint`: FNV-1a 64 of every output of the warm-up batch. At `--batch-log2 13` every implementation so far + prints `f99fb375b3abeaf5` for igneum-genesis-mh (Apple OpenCL, pocl on both exchange paths, the emulator in seven + configurations); an AMD run at `--batch-log2 13` must print the same value. At the default 2^24 the value is a new + reference to compare AMD against the next machine. +8. `timed:` with device and wall time, then the `rate` line: Mhash/s and GB/s useful. +9. The summary table and `OVERALL: PASS`. Exit code 0 on PASS, 1 on FAIL, 2 on an OpenCL error or a build failure (the + build log is printed in full). + +A FAIL with a few lanes differing points at the exchange; a FAIL in every lane at an arithmetic op or the dataset; a +cache FAIL at the ChaCha fill. Send the whole printout either way, including the `--list` output and, on Linux, +`clinfo` and the ROCm or Mesa version, or on Windows the Adrenalin driver version from the AMD Software panel. + +## What to paste back + +All seven text files above, unedited. From them the bench log gets: device name and driver string, the `exchange:` +line, the cache check line, the 96/96 vector result per pack, the rate at 1 GiB, and the sweep table. If a run looks +odd, run it again and keep both. + +## Results so far (no AMD) + +| Machine, runtime | Exchange | igneum-genesis-mh | closed-form packs | Rate at 1 GiB | +|---|---|---|---|---| +| Apple M5 Max, Apple OpenCL 1.2 (deprecated runtime) | local memory (runtime lists no sub-group extension) | cache FNV = Mac, 96/96; also 96/96 at `--group-warps 2` and `4` | 96/96 and 96/96 | 45.03 Mhash/s, 18.73 GB/s useful (Metal on the same chip: 45.2). Apple OpenCL on the M5 Max, not an AMD number | +| Apple M5 Max, pocl 7.2 CPU device, OpenCL 3.0, LLVM 23 | `sub_group_shuffle_xor` (auto and `--exchange subgroup`; pocl's `clGetKernelSubGroupInfoKHR` fails with CL_INVALID_OPERATION, the probe kernel reports 32), and `--exchange local` | cache FNV = Mac, 96/96 on both paths, same batch fingerprint as Apple and the emulator at 2^13 | not run | CPU, not meaningful | +| CPU emulator (`emu/`), 7 configurations incl. sub-group 64 | both | 96/96 in every configuration, identical batch fingerprint `f99fb375b3abeaf5` | not run | none | + +Sweep on Apple OpenCL (3 batches, wall time): 4 MiB 573.7, 64 MiB 178.9, 256 MiB 94.3, 512 MiB 68.8, 1024 MiB 45.0 Mhash/s. +Full tables in `docs/bench-log.md`. + +## Checking a pack without a GPU + +``` +emu/emu.sh igneum-genesis-mh 0 32 --sg 32 # local-memory exchange, work-group 32, 32-wide sub-group +emu/emu.sh igneum-genesis-mh 0 64 --sg 64 # local-memory exchange on a wave64 model, two units per work-group +emu/emu.sh igneum-genesis-mh 1 32 --sg 32 # sub_group_shuffle_xor, 32-wide sub-group +emu/emu.sh igneum-genesis-mh 1 64 --sg 64 # sub_group_shuffle_xor over a 64-wide sub-group (two units in one wave) +``` + +The emulator compiles `kernel.cl` unchanged as C++ (`emu/emu_opencl.h` supplies the OpenCL built-ins; the kernel's own +prelude includes it when no OpenCL compiler is present) and runs work-groups on host threads with a barrier inside +every exchange. It prints the cache check, dataset self-test, vectors and a fingerprint of the whole 2^13-hash batch, +so configurations can be compared for every nonce, not only the vector warps. It says nothing about any GPU. + +## Regenerating a pack + +``` +cd proto-metal +swiftc -O -o igneum-bench main.swift -framework Metal +./igneum-bench --seed igneum-genesis --export-pack ../proto-cuda/packs/igneum-genesis-mh +./igneum-bench --closed-form --seed igneum-genesis --export-pack ../proto-cuda/packs/igneum-genesis +./igneum-bench --closed-form --seed igneum-hourly --export-pack ../proto-cuda/packs/igneum-hourly +``` + +The exporter writes `kernel.cl` next to `kernel.cu` from the same instruction list, refuses to write unless the Metal +GPU matches the CPU interpreter on all 96 vector outputs, and for memory-hard packs unless the GPU cache equals the +CPU cache word for word. The memory-hard core is emitted once (`emitMemhardCore`) in three dialects (Metal, CUDA C++, +OpenCL C), so the constants in `kernel.cl` are the literals of `memhard.h`. diff --git a/proto-opencl/WAVEFRONT.md b/proto-opencl/WAVEFRONT.md new file mode 100644 index 000000000..6bfb2e8e7 --- /dev/null +++ b/proto-opencl/WAVEFRONT.md @@ -0,0 +1,98 @@ +# Wavefront width and the 32-lane verification unit + +Written 3 October 2026 for the OpenCL path (`proto-opencl`). It records the one place where AMD hardware differs +from Apple and NVIDIA in a way the lottery hash can feel, and how `kernel.cl` is built so that it cannot. + +## The unit is 32 lanes, by definition + +The program pack defines a hash over 32 lanes. Six of its 64 instructions are exchanges (`shfl`): lane `l` reads a +register from lane `l ^ m` with `m` in {1, 2, 4, 8, 16}. The CPU verifier (`proto-metal/main.swift`, `cpuWarp`) +models exactly 32 lanes. Metal runs it on a 32-wide SIMD group (`simd_shuffle_xor`), CUDA on a 32-wide warp +(`__shfl_xor_sync`). The vectors in every pack are 3 warps x 32 lanes x 64 bits. Nothing in the definition mentions +hardware; 32 is a parameter of the hash. + +## What AMD hardware does + +| Architecture | Native wave width | Notes | +|---|---|---| +| GCN (Radeon HD 7000 to Vega, approximate) | 64 | always wave64 | +| CDNA (Instinct MI100 to MI300, approximate) | 64 | always wave64 | +| RDNA 1 to 4 (RX 5000 to RX 9000, approximate) | 32 or 64 | the shader compiler picks per kernel; compute kernels are often wave64 unless the compiler decides otherwise, and nothing in OpenCL lets the program demand wave32 | + +Figures from memory, labelled approximate. NVIDIA is 32 everywhere; Apple is 32 on every chip the Mac tool has +seen (it warns if `threadExecutionWidth` is not 32). In OpenCL terms the hardware wave is the sub-group: +`get_sub_group_size()` and `clGetKernelSubGroupInfoKHR(CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE_KHR)` report it. + +So on an AMD card a sub-group may be 64 lanes, and a sub-group shuffle (`sub_group_shuffle_xor`) spans 64 lanes. +Two things could go wrong if the kernel simply used sub-group shuffles: + +1. A wave64 sub-group holds two logical 32-lane units. `lane ^ m` with `m < 32` never leaves the unit's own aligned + half, so the xor exchange itself would still be correct. The emulator below demonstrates this. But a + `sub_group_broadcast(x, 0)` (used only by the optional wide-load lever `wload`, off by default) would broadcast + lane 0 of the wave, which is wrong for the upper unit. +2. The mapping from sub-group lane to work-item is not something the OpenCL specification promises to be linear + and aligned. In practice it is (consecutive work-items in linear local id order), on AMD, NVIDIA, Intel and + pocl alike, but "in practice" is not a guarantee a consensus rule should rest on. + +## The rule kernel.cl follows + +The exchange is selected at compile time by `IGNEUM_EXCHANGE`, and `proto-opencl/host.c` chooses the value: + +| `IGNEUM_EXCHANGE` | Exchange | When host.c selects it | +|---|---|---| +| 0 (default) | local memory: each lane writes its register to `__local` memory, `barrier(CLK_LOCAL_MEM_FENCE)`, reads slot `lid ^ m` | always available; chosen whenever any condition below fails | +| 1 | `sub_group_shuffle_xor` (extension `cl_khr_subgroup_shuffle`, OpenCL C 2.0 or 3.0) | the device lists the extension, `--group-warps 1` (work-group of exactly 32, so one work-group is one unit), the variant compiles, and `clGetKernelSubGroupInfoKHR` on `igneum_hash` reports a sub-group size of exactly 32 for a 32-item work-group | +| 2 | `intel_sub_group_shuffle_xor` (`cl_intel_subgroups`) | as 1, on Intel devices without the khr shuffle extension | + +In words: the 32-lane group is defined by the work-group of 32 (one work-group = one unit) whenever sub-group +shuffles are used, and the local-memory exchange is used whenever the sub-group size is not exactly 32. On a +wave64 device the queried size is 64, so AMD GCN, CDNA and RDNA-in-wave64 all take the local-memory path, and the +hash does not depend on the wave width at all: a work-group barrier and `__local` memory mean the same thing at any +wave width. If the sub-group size cannot be queried (`clGetKernelSubGroupInfoKHR` missing, as on pocl), host.c +runs a probe kernel from the same build that reports `get_sub_group_size()` for a 32-item work-group; this is +weaker than the per-kernel query because a compiler may choose the wave width per kernel, and the printout says so. +`--exchange local` forces path 0 on any device; `--exchange subgroup` demands path 1 or 2 and fails otherwise. +The run prints which path it took and why, on the `exchange:` line. + +The local-memory exchange uses two buffers of `IGNEUM_GROUP` words that alternate (counter `xk`), so one barrier per +exchange suffices: a lane can only overwrite buffer `b` at exchange `k + 2` after passing barrier `k + 1`, and every +lane reaches barrier `k + 1` only after its read of buffer `b` at exchange `k`. Control flow is uniform (the program +has no branches), so every work-item reaches every barrier. `--group-warps W` packs W units into one work-group of +32 W items; `lid ^ m` stays inside the unit because `m < 32`, and the Apple runs below show W = 1, 2 and 4 bit-exact. + +## What was proven without AMD silicon (3 October 2026) + +| Check | Path | Result | +|---|---|---| +| Apple M5 Max, Apple OpenCL 1.2 runtime, pack igneum-genesis-mh | local memory (the runtime lists no sub-group extension), work-group 32 | cache FNV `48c4f5bf24166b2e` = Mac, 96/96 vectors standalone and in batch, 45.03 Mhash/s at 1 GiB | +| Same, `--exchange local --group-warps 2` and `--group-warps 4` | local memory, work-groups of 64 and 128 | 96/96 and 96/96 | +| Same, closed-form packs igneum-genesis and igneum-hourly | local memory | 96/96 and 96/96 | +| pocl 7.2 CPU device (OpenCL 3.0, LLVM 23), Khronos ICD loader, `--exchange auto` and `--exchange subgroup` | `sub_group_shuffle_xor`, OpenCL C 3.0 (pocl's `clGetKernelSubGroupInfoKHR` fails with CL_INVALID_OPERATION, the probe kernel reports a sub-group of 32 for a 32-item work-group) | cache FNV = Mac, 96/96 | +| Same, `--exchange local` | local memory | cache FNV = Mac, 96/96 | +| CPU emulator, `IGNEUM_EXCHANGE 0`, work-group 32, sub-group 32 | local memory | PASS, batch fingerprint `f99fb375b3abeaf5` | +| CPU emulator, `IGNEUM_EXCHANGE 0`, work-group 64, sub-group 64 (wave64, two units per wave) | local memory | PASS, same fingerprint | +| CPU emulator, `IGNEUM_EXCHANGE 0`, work-group 32, sub-group 64 (wave64 half empty) | local memory | PASS, same fingerprint | +| CPU emulator, `IGNEUM_EXCHANGE 1`, work-group 32, sub-group 32 | `sub_group_shuffle_xor` | PASS, same fingerprint | +| CPU emulator, `IGNEUM_EXCHANGE 1`, work-group 32, sub-group 64 | `sub_group_shuffle_xor` over a half-empty wave64 | PASS, same fingerprint | +| CPU emulator, `IGNEUM_EXCHANGE 1`, work-group 64, sub-group 64 (wave64, two units in one shuffle domain) | `sub_group_shuffle_xor` | PASS, same fingerprint | +| CPU emulator, `IGNEUM_EXCHANGE 1`, work-group 64, sub-group 32 (two sub-groups per work-group) | `sub_group_shuffle_xor` | PASS, same fingerprint | + +The fingerprint is FNV-1a 64 over the 2^13 outputs of the batch at base nonce 0; seven emulator configurations +produced the same one, and host.c prints the same fingerprint for its warm-up batch (`--batch-log2 13`): Apple's +OpenCL (local memory) and pocl (sub-group shuffles, and local memory) all print `f99fb375b3abeaf5`. So both exchange +implementations, both wave widths, two real OpenCL compilers and the emulator give identical hashes for every nonce +in the batch, not only for the three vector warps. An AMD run at `--batch-log2 13` must print the same value. The emulator rows with sub-group 64 and `IGNEUM_EXCHANGE 1` show that the xor +exchange alone would survive a wave64 sub-group (point 1 above); host.c still refuses that configuration on a real +device because of point 2 and because of `sub_group_broadcast`. The conservative rule costs nothing in correctness +and, on a wave64 card, the local-memory path is what runs. + +## What is not proven here + +- No AMD compiler has compiled `kernel.cl`, and no AMD device has run it. The emulator is clang; pocl is LLVM on a + CPU; Apple's OpenCL is Apple's compiler. Syntax every one of them accepts can still trip AMD's front end, which + would be a build log (printed in full by host.c), not a silent difference. +- No AMD hash rate exists. The cost of the local-memory exchange relative to sub-group shuffles on AMD is unknown; + on Apple the local-memory path runs at the same 45 Mhash/s as Metal's `simd_shuffle_xor` kernel because the hash is + bound by the 104 random dataset loads, and the same is expected elsewhere, but expected is not measured. +- Whether RDNA compiles `igneum_hash` as wave32 (sub-group 32, path 1) or wave64 (path 0) is a driver decision. The + run will say which on the `exchange:` line. Both are correct by construction; only the speed may differ. diff --git a/proto-opencl/build.bat b/proto-opencl/build.bat new file mode 100644 index 000000000..1fa9ac02b --- /dev/null +++ b/proto-opencl/build.bat @@ -0,0 +1,59 @@ +@echo off +rem Build igneum-bench-cl for one program pack (Windows, MSVC). +rem Run from an "x64 Native Tools Command Prompt for VS 2022" (or 2026 with any toolset: this is plain C, no nvcc involved). +rem Needs the OpenCL headers (CL\cl.h) and the ICD loader import library (OpenCL.lib). The script looks, in order, at: +rem OPENCL_SDK a Khronos OpenCL-SDK or vcpkg tree: %OPENCL_SDK%\include\CL\cl.h and %OPENCL_SDK%\lib\OpenCL.lib +rem CUDA_PATH the CUDA Toolkit ships both: %CUDA_PATH%\include\CL\cl.h and %CUDA_PATH%\lib\x64\OpenCL.lib +rem OCL_ROOT the old AMD APP SDK layout: %OCL_ROOT%\include and %OCL_ROOT%\lib\x86_64 +rem The runtime itself (OpenCL.dll in System32 plus the vendor ICD) comes with the AMD Adrenalin driver; no SDK needed to run. +rem Usage: build.bat [pack] +rem pack directory name under ..\proto-cuda\packs\ (default igneum-genesis-mh) +setlocal +cd /d "%~dp0" + +set PACK=%1 +if "%PACK%"=="" set PACK=igneum-genesis-mh +set PACKDIR=..\proto-cuda\packs\%PACK% + +where cl >nul 2>nul +if errorlevel 1 ( + echo cl.exe not found. Open an "x64 Native Tools Command Prompt" ^(Visual Studio Build Tools, Desktop development with C++^) and run this again. + exit /b 1 +) +if not exist "%PACKDIR%\kernel.cl" ( + echo no pack at %PACKDIR% ^(expected kernel.cl, program.h, vectors.h^) + exit /b 1 +) + +set INC= +set LIB= +if not "%OPENCL_SDK%"=="" if exist "%OPENCL_SDK%\include\CL\cl.h" ( + set INC=%OPENCL_SDK%\include + if exist "%OPENCL_SDK%\lib\OpenCL.lib" set LIB=%OPENCL_SDK%\lib\OpenCL.lib + if exist "%OPENCL_SDK%\lib\x64\OpenCL.lib" set LIB=%OPENCL_SDK%\lib\x64\OpenCL.lib +) +if "%LIB%"=="" if not "%CUDA_PATH%"=="" if exist "%CUDA_PATH%\include\CL\cl.h" ( + set INC=%CUDA_PATH%\include + set LIB=%CUDA_PATH%\lib\x64\OpenCL.lib +) +if "%LIB%"=="" if not "%OCL_ROOT%"=="" if exist "%OCL_ROOT%\include\CL\cl.h" ( + set INC=%OCL_ROOT%\include + set LIB=%OCL_ROOT%\lib\x86_64\OpenCL.lib +) +if "%LIB%"=="" ( + echo No OpenCL SDK found. Set OPENCL_SDK to a tree with include\CL\cl.h and lib\OpenCL.lib, for example: + echo vcpkg install opencl:x64-windows then set OPENCL_SDK=C:\vcpkg\installed\x64-windows + echo or the Khronos OpenCL-SDK release zip, or install the CUDA Toolkit ^(it ships CL\cl.h and OpenCL.lib^). + exit /b 1 +) + +echo using headers %INC% and import library %LIB% +echo cl /nologo /O2 /W3 /std:c11 /I "%INC%" /I "%PACKDIR%" /DIGNEUM_KERNEL_PATH="\"%PACKDIR:\=/%/kernel.cl\"" /Fe:igneum-bench-cl-%PACK%.exe host.c /link "%LIB%" +cl /nologo /O2 /W3 /std:c11 /I "%INC%" /I "%PACKDIR%" /DIGNEUM_KERNEL_PATH="\"%PACKDIR:\=/%/kernel.cl\"" /Fe:igneum-bench-cl-%PACK%.exe host.c /link "%LIB%" +if errorlevel 1 ( + echo build failed. If cl.exe rejects /std:c11, drop that flag: the code is C99 and MSVC's default C mode accepts it. + exit /b 1 +) +del host.obj >nul 2>nul +echo built igneum-bench-cl-%PACK%.exe +endlocal diff --git a/proto-opencl/build.sh b/proto-opencl/build.sh new file mode 100755 index 000000000..8c118a3aa --- /dev/null +++ b/proto-opencl/build.sh @@ -0,0 +1,48 @@ +#!/usr/bin/env bash +# Build igneum-bench-cl for one program pack (macOS or Linux). +# Usage: ./build.sh [pack] [loader] +# pack directory name under ../proto-cuda/packs/ (default igneum-genesis-mh) +# loader macOS only: "apple" (default) links Apple's deprecated OpenCL.framework; "khr" links the Khronos ICD loader +# from Homebrew (brew install opencl-headers opencl-icd-loader pocl) so a CPU OpenCL 3.0 device can be tested. +# On Linux the argument is ignored and -lOpenCL is used (AMD ROCm, AMD Adrenalin for Linux, Mesa rusticl/clover, +# NVIDIA or Intel runtimes all register with the system ICD loader). +set -euo pipefail +cd "$(dirname "$0")" + +PACK="${1:-igneum-genesis-mh}" +LOADER="${2:-apple}" +PACKDIR="../proto-cuda/packs/$PACK" + +if [ ! -f "$PACKDIR/kernel.cl" ]; then + echo "no pack at $PACKDIR (expected kernel.cl, program.h, vectors.h); regenerate with proto-metal/igneum-bench --export-pack" >&2 + exit 1 +fi + +CC="${CC:-cc}" +COMMON=(-std=c99 -O2 -Wall -Wextra -I "$PACKDIR" -DIGNEUM_KERNEL_PATH="\"$PACKDIR/kernel.cl\"") +OUT="igneum-bench-cl-$PACK" + +case "$(uname -s)" in + Darwin) + if [ "$LOADER" = "khr" ]; then + PREFIX="$(brew --prefix 2>/dev/null || echo /opt/homebrew)" + OUT="$OUT-khr" + set -x + "$CC" "${COMMON[@]}" -DIGNEUM_KHR_HEADERS -I "$PREFIX/opt/opencl-headers/include" -o "$OUT" host.c \ + -L "$PREFIX/opt/opencl-icd-loader/lib" -lOpenCL + set +x + echo "built ./$OUT (Khronos ICD loader; run with OCL_ICD_VENDORS=$PREFIX/etc/OpenCL/vendors)" + else + set -x + "$CC" "${COMMON[@]}" -Wno-deprecated-declarations -o "$OUT" host.c -framework OpenCL + set +x + echo "built ./$OUT (Apple OpenCL 1.2, deprecated since macOS 10.14; correctness check only, not an AMD number)" + fi + ;; + *) + set -x + "$CC" "${COMMON[@]}" -o "$OUT" host.c -lOpenCL -ldl + set +x + echo "built ./$OUT" + ;; +esac diff --git a/proto-opencl/emu/emu.sh b/proto-opencl/emu/emu.sh new file mode 100755 index 000000000..e5d49f36d --- /dev/null +++ b/proto-opencl/emu/emu.sh @@ -0,0 +1,27 @@ +#!/usr/bin/env bash +# CPU emulation of the OpenCL pack, for checking kernel.cl on a machine WITHOUT an AMD GPU (any clang++ or g++). +# It compiles the generated kernel.cl as C++ against emu_opencl.h and runs the kernels on host threads with a chosen +# sub-group width (32, or 64 to model a wave64 GPU running two logical 32-lane units in one wave). +# Usage: emu/emu.sh [exchange] [group] [emulator args...] +# pack directory name under ../proto-cuda/packs/ (default igneum-genesis-mh) +# exchange IGNEUM_EXCHANGE: 0 = local-memory exchange, 1 = sub_group_shuffle_xor (default 0) +# group IGNEUM_GROUP: work-group size of igneum_hash, 32 or 64 (default 32) +# rest passed to the emulator: --sg 32|64 (default 32), --batch-log2 B (default 13), --dataset-mib N +# Example (wave64, two units per wave, sub-group shuffles): emu/emu.sh igneum-genesis-mh 1 64 --sg 64 +# Only PASS/FAIL matters. Rates are not printed because they would be meaningless. +set -euo pipefail +HERE="$(cd "$(dirname "$0")" && pwd)" +PACK="${1:-igneum-genesis-mh}"; shift || true +EXCHANGE="${1:-0}"; shift || true +GROUP="${1:-32}"; shift || true +PACKDIR="$(cd "$HERE/../../proto-cuda/packs/$PACK" && pwd)" +BUILD="$HERE/build-$PACK-e$EXCHANGE-g$GROUP" +mkdir -p "$BUILD" +CXX="${CXX:-c++}" +FLAGS=(-std=c++17 -O2 -Wall -Wextra -I "$HERE" -I "$PACKDIR" -D "IGNEUM_GROUP=$GROUP" -D "IGNEUM_EXCHANGE=$EXCHANGE") +# kernel.cl is compiled as C++ unchanged: its prelude includes emu_opencl.h when __OPENCL_VERSION__ is absent. +"$CXX" "${FLAGS[@]}" -x c++ -c -o "$BUILD/kernel_cl.o" "$PACKDIR/kernel.cl" +"$CXX" "${FLAGS[@]}" -c -o "$BUILD/emu_main.o" "$HERE/emu_main.cpp" +"$CXX" -o "$BUILD/igneum-emu-cl" "$BUILD/kernel_cl.o" "$BUILD/emu_main.o" -pthread +echo "compiled $BUILD/igneum-emu-cl (CPU emulation of kernel.cl, not a GPU build)" +exec "$BUILD/igneum-emu-cl" "$@" diff --git a/proto-opencl/emu/emu_main.cpp b/proto-opencl/emu/emu_main.cpp new file mode 100644 index 000000000..532718faf --- /dev/null +++ b/proto-opencl/emu/emu_main.cpp @@ -0,0 +1,301 @@ +// CPU emulator driver for the OpenCL pack: runs the generated kernel.cl (compiled as C++ through emu_opencl.h) on +// host threads with a configurable sub-group width, and checks cache, dataset and vectors against the pack. +// +// Model. A launch of global G work-items with work-group size L spawns L host threads; thread l plays work-item l +// of every work-group in turn (group 0, 1, 2, ...), exactly like proto-cuda/emu. barrier() synchronises the L +// threads of the current group. Sub-groups are consecutive runs of SG work-items inside the group (SG = --sg 32 or +// 64; the last run is clipped to the group size); sub_group_shuffle_xor exchanges through a slot array with a +// barrier among the sub-group's threads, so a 64-wide sub-group really carries two logical 32-lane units. +// Local memory is one arena per work-group instance, never shared between groups, as on hardware. +// +// Compile-time (as on the device): IGNEUM_GROUP (work-group size of igneum_hash) and IGNEUM_EXCHANGE (0 local memory, +// 1 sub-group shuffles). Runtime: --sg 32|64, --batch-log2 B (default 13), --dataset-mib N (default the pack size). +// Only PASS/FAIL matters here. Rates are meaningless and not printed. +#include "emu_opencl.h" +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#define IGNEUM_NO_CUDA +#include "program.h" +#include "vectors.h" +#ifndef IGNEUM_DATASET_MODE +#define IGNEUM_DATASET_MODE 0 +#endif +#if IGNEUM_DATASET_MODE == 1 +#include "memhard.h" +#endif +#ifndef IGNEUM_GROUP +#define IGNEUM_GROUP 32 +#endif +#ifndef IGNEUM_EXCHANGE +#define IGNEUM_EXCHANGE 0 +#endif + +// Kernels from kernel.cl (compiled as a separate C++ translation unit with the same defines). +void igneum_hash(const uint* ds, ulong* out, uint baseNonce, uint mask); +#if IGNEUM_DATASET_MODE == 1 +void igneum_cache_fill(uint* cache, uint nSegments); +void igneum_build(uint* ds, const uint* cache, uint nItems); +#else +void igneum_fill(uint* ds, uint n, uint d0, uint d1); +#endif + +// --------------------------------------------------------------------------------------------- +// Runtime + +struct Barrier { + std::mutex m; + std::condition_variable cv; + unsigned size = 0, arrived = 0, generation = 0; + void wait() { + std::unique_lock lk(m); + unsigned gen = generation; + if (++arrived == size) { arrived = 0; ++generation; cv.notify_all(); } + else cv.wait(lk, [&] { return gen != generation; }); + } +}; + +struct SubGroup { + Barrier bar; + uint slot[64]; + unsigned size = 0; +}; + +struct Launch { + unsigned local = 0, groups = 0, sg = 32; + Barrier groupBar; + std::vector> subs; + std::mutex arenaMutex; + std::unordered_map> arenas; // group index -> local memory words + unsigned arenaWords = 0; +}; + +static thread_local Launch* tlLaunch = nullptr; +static thread_local unsigned tlLid = 0, tlGroup = 0; +static unsigned gSubGroupWidth = 32; + +size_t get_global_id(uint) { return (size_t)tlGroup * tlLaunch->local + tlLid; } +size_t get_local_id(uint) { return tlLid; } +size_t get_group_id(uint) { return tlGroup; } +size_t get_local_size(uint) { return tlLaunch->local; } +size_t get_global_size(uint) { return (size_t)tlLaunch->local * tlLaunch->groups; } +uint get_sub_group_size(void) { return tlLaunch->subs[tlLid / tlLaunch->sg]->size; } +uint get_sub_group_local_id(void) { return tlLid % tlLaunch->sg; } +uint get_sub_group_id(void) { return tlLid / tlLaunch->sg; } +uint get_num_sub_groups(void) { return (uint)tlLaunch->subs.size(); } + +void barrier(int) { tlLaunch->groupBar.wait(); } + +uint sub_group_shuffle_xor(uint v, uint mask) { + SubGroup* s = tlLaunch->subs[tlLid / tlLaunch->sg].get(); + unsigned lane = tlLid % tlLaunch->sg; + unsigned partner = lane ^ mask; + if (partner >= s->size) { + std::fprintf(stderr, "emu: sub_group_shuffle_xor partner lane %u outside the sub-group of %u lanes (undefined on hardware)\n", partner, s->size); + std::exit(3); + } + s->slot[lane] = v; + s->bar.wait(); + uint r = s->slot[partner]; + s->bar.wait(); + return r; +} +uint intel_sub_group_shuffle_xor(uint v, uint mask) { return sub_group_shuffle_xor(v, mask); } + +uint sub_group_broadcast(uint v, uint laneSrc) { + SubGroup* s = tlLaunch->subs[tlLid / tlLaunch->sg].get(); + unsigned lane = tlLid % tlLaunch->sg; + s->slot[lane] = v; + s->bar.wait(); + uint r = s->slot[laneSrc % s->size]; + s->bar.wait(); + return r; +} + +uint* emu_local_words(uint n) { + Launch* L = tlLaunch; + std::lock_guard lk(L->arenaMutex); + auto it = L->arenas.find(tlGroup); + if (it == L->arenas.end()) { + std::unique_ptr a(new uint[n]()); + it = L->arenas.emplace(tlGroup, std::move(a)).first; + L->arenaWords = n; + } + return it->second.get(); +} + +template static void emu_launch(F f, size_t global, unsigned local, A... args) { + if (global % local != 0) { std::fprintf(stderr, "emu: global %zu not a multiple of local %u\n", global, local); std::exit(3); } + Launch L; + L.local = local; L.groups = (unsigned)(global / local); L.sg = gSubGroupWidth; + L.groupBar.size = local; + unsigned nSub = (local + L.sg - 1) / L.sg; + for (unsigned k = 0; k < nSub; ++k) { + L.subs.emplace_back(new SubGroup()); + L.subs.back()->size = (local - k * L.sg) < L.sg ? (local - k * L.sg) : L.sg; + L.subs.back()->bar.size = L.subs.back()->size; + } + std::vector ts; + for (unsigned t = 0; t < local; ++t) { + ts.emplace_back([&, t]() { + tlLaunch = &L; tlLid = t; + for (unsigned g = 0; g < L.groups; ++g) { tlGroup = g; f(args...); } + }); + } + for (auto& th : ts) th.join(); +} + +// --------------------------------------------------------------------------------------------- +// Checks + +static uint64_t fnv1a64(const void* p, size_t n) { + const uint8_t* b = (const uint8_t*)p; + uint64_t h = 0xcbf29ce484222325ull; + for (size_t i = 0; i < n; ++i) { h ^= b[i]; h *= 0x100000001b3ull; } + return h; +} + +#if IGNEUM_DATASET_MODE == 0 +static uint32_t host_ds_elem(uint32_t i, uint32_t d0, uint32_t d1) { + uint32_t x = i ^ d0; + x *= 0x9E3779B1u; x ^= x >> 15; + x += d1; + x *= 0x85EBCA77u; x ^= x >> 13; + x *= 0xC2B2AE3Du; x ^= x >> 16; + return x; +} +#endif + +static bool compareWarp(const uint64_t* got, const uint64_t* want, uint32_t base, const char* how) { + int bad = 0, first = -1; + for (int l = 0; l < 32; ++l) if (got[l] != want[l]) { if (first < 0) first = l; ++bad; } + if (bad == 0) std::printf("verify warp base %u %s: PASS\n", base, how); + else std::printf("verify warp base %u %s: FAIL %d of 32 lanes differ, first lane %d: emu=%016llx expected=%016llx\n", + base, how, bad, first, (unsigned long long)got[first], (unsigned long long)want[first]); + return bad == 0; +} + +static double wallMs() { + using namespace std::chrono; + return duration(steady_clock::now().time_since_epoch()).count(); +} + +int main(int argc, char** argv) { + int batchLog2 = 13; + int datasetMib = (int)(((1ull << IGNEUM_DATASET_LOG2) * 4ull) >> 20); + for (int i = 1; i < argc; ++i) { + std::string a = argv[i]; + if (a == "--sg" && i + 1 < argc) gSubGroupWidth = (unsigned)std::atoi(argv[++i]); + else if (a == "--batch-log2" && i + 1 < argc) batchLog2 = std::atoi(argv[++i]); + else if (a == "--dataset-mib" && i + 1 < argc) datasetMib = std::atoi(argv[++i]); + else { std::printf("usage: igneum-emu-cl [--sg 32|64] [--batch-log2 13] [--dataset-mib N]\n"); return 2; } + } + if (gSubGroupWidth != 32 && gSubGroupWidth != 64) { std::printf("--sg must be 32 or 64\n"); return 2; } + if (IGNEUM_EXCHANGE != 0 && IGNEUM_GROUP != 32 && IGNEUM_GROUP != 64) { std::printf("sub-group exchange needs IGNEUM_GROUP 32 or 64 here\n"); return 2; } + + const uint32_t words = (uint32_t)(((uint64_t)datasetMib << 20) / 4ull); + const uint32_t mask = words - 1u; + const bool atPackSize = (words == (1u << IGNEUM_DATASET_LOG2)); + std::printf("igneum-emu-cl pack \"%s\" CPU EMULATION of kernel.cl (not a GPU; PASS/FAIL only)\n", IGNEUM_SEED_STRING); + std::printf("configuration: IGNEUM_GROUP %d, IGNEUM_EXCHANGE %d (%s), emulated sub-group width %u%s, dataset %d MiB, batch 2^%d\n", + IGNEUM_GROUP, IGNEUM_EXCHANGE, IGNEUM_EXCHANGE == 0 ? "local-memory exchange with barrier" : "sub_group_shuffle_xor", + gSubGroupWidth, gSubGroupWidth == 64 ? " (two logical 32-lane units per wave)" : "", datasetMib, batchLog2); + std::printf("hardware threads: %u\n", std::thread::hardware_concurrency()); + + bool overall = true; + std::vector ds(words); +#if IGNEUM_DATASET_MODE == 1 + const uint32_t cacheWords = 1u << IGNEUM_CACHE_LOG2_WORDS; + std::vector cache(cacheWords), hostCache(cacheWords); + double t0 = wallMs(); + emu_launch(igneum_cache_fill, (size_t)IGNEUM_CACHE_SEGMENTS, 256u, cache.data(), (uint)IGNEUM_CACHE_SEGMENTS); + double t1 = wallMs(); + for (uint32_t seg = 0; seg < IGNEUM_CACHE_SEGMENTS; ++seg) mh_cache_segment(hostCache.data(), seg); + double t2 = wallMs(); + bool same = std::memcmp(cache.data(), hostCache.data(), (size_t)cacheWords * 4u) == 0; + uint64_t fnv = fnv1a64(cache.data(), (size_t)cacheWords * 4u); + bool fnvOk = fnv == IGNEUM_CACHE_FNV64; + bool headOk = std::memcmp(cache.data(), IGNEUM_CACHE_HEAD, 64) == 0; + bool lastOk = std::memcmp(cache.data() + cacheWords - 16u, IGNEUM_CACHE_LAST, 64) == 0; + std::printf("cache: emulated igneum_cache_fill %.0f ms (256 threads), host memhard.h one thread %.0f ms\n", t1 - t0, t2 - t1); + std::printf("cache check: %s (emulated kernel == host all %u words %s, FNV-1a 64 %016llx vs Mac %016llx %s, head %s, last line %s)\n", + (same && fnvOk && headOk && lastOk) ? "PASS" : "FAIL", cacheWords, same ? "PASS" : "FAIL", + (unsigned long long)fnv, (unsigned long long)IGNEUM_CACHE_FNV64, fnvOk ? "PASS" : "FAIL", headOk ? "PASS" : "FAIL", lastOk ? "PASS" : "FAIL"); + overall = overall && same && fnvOk && headOk && lastOk; + t0 = wallMs(); + emu_launch(igneum_build, (size_t)(words / 16u), 256u, ds.data(), (const uint*)cache.data(), (uint)(words / 16u)); + std::printf("dataset: emulated igneum_build %.0f ms for 2^%u items\n", wallMs() - t0, (unsigned)(IGNEUM_DATASET_LOG2 - 4)); +#else + double t0 = wallMs(); + emu_launch(igneum_fill, (size_t)words, 256u, ds.data(), words, (uint)IGNEUM_DAY0, (uint)IGNEUM_DAY1); + std::printf("dataset: emulated igneum_fill %.0f ms\n", wallMs() - t0); +#endif + + // Dataset self-test, same shape as host.c. + { + int badHead = 0, badRnd = 0, badSample = 0, nSample = 0; + bool lastOk = true; + for (int i = 0; i < 16; ++i) if (ds[i] != IGNEUM_DS_HEAD[i]) ++badHead; + if (atPackSize) lastOk = (ds[IGNEUM_DS_LAST_INDEX] == IGNEUM_DS_LAST); + uint64_t s = 0x9E3779B97F4A7C15ull ^ (uint64_t)words; + for (int k = 0; k < 64; ++k) { + s += 0x9E3779B97F4A7C15ull; + uint64_t z = s; + z = (z ^ (z >> 30)) * 0xBF58476D1CE4E5B9ull; + z = (z ^ (z >> 27)) * 0x94D049BB133111EBull; + z ^= z >> 31; + uint32_t idx = (uint32_t)z & mask; +#if IGNEUM_DATASET_MODE == 1 + uint32_t want = mh_word(hostCache.data(), idx); +#else + uint32_t want = host_ds_elem(idx, IGNEUM_DAY0, IGNEUM_DAY1); +#endif + if (ds[idx] != want) ++badRnd; + } +#ifdef IGNEUM_DS_SAMPLES + for (int k = 0; k < IGNEUM_DS_SAMPLES; ++k) { + if (IGNEUM_DS_SAMPLE_INDEX[k] > mask) continue; + ++nSample; + if (ds[IGNEUM_DS_SAMPLE_INDEX[k]] != IGNEUM_DS_SAMPLE_VALUE[k]) ++badSample; + } +#endif + bool dsPass = badHead == 0 && lastOk && badRnd == 0 && badSample == 0; + std::printf("dataset self-test: %s (head 16 vs Mac %s, element [MASK] %s, 64 random points vs host %s, %d Mac samples %s)\n", + dsPass ? "PASS" : "FAIL", badHead == 0 ? "PASS" : "FAIL", atPackSize ? (lastOk ? "PASS" : "FAIL") : "skipped", + badRnd == 0 ? "PASS" : "FAIL", nSample, badSample == 0 ? "PASS" : "FAIL"); + overall = overall && dsPass; + } + + if (!atPackSize) { std::printf("vectors: skipped (not the pack size)\nOVERALL: %s\n", overall ? "PASS" : "FAIL"); return overall ? 0 : 1; } + + // Vectors standalone: one work-group of IGNEUM_GROUP items per base nonce (the first 32 are the vector warp). + const uint32_t nonces = 1u << batchLog2; + std::vector out(nonces); + for (int w = 0; w < IGNEUM_VEC_WARPS; ++w) { + emu_launch(igneum_hash, (size_t)IGNEUM_GROUP, (unsigned)IGNEUM_GROUP, (const uint*)ds.data(), out.data(), (uint)IGNEUM_VEC_BASE[w], mask); + char how[96]; + std::snprintf(how, sizeof(how), "standalone, work-group %d, sub-group width %u", IGNEUM_GROUP, gSubGroupWidth); + overall = compareWarp((const uint64_t*)out.data(), IGNEUM_VEC_OUT[w], IGNEUM_VEC_BASE[w], how) && overall; + } + // In batch: every vector warp that fits in 2^batchLog2 nonces. + emu_launch(igneum_hash, (size_t)nonces, (unsigned)IGNEUM_GROUP, (const uint*)ds.data(), out.data(), 0u, mask); + for (int w = 0; w < IGNEUM_VEC_WARPS; ++w) { + if ((uint64_t)IGNEUM_VEC_BASE[w] + 32ull > nonces) { std::printf("verify warp base %u in batch: skipped (batch has %u nonces)\n", IGNEUM_VEC_BASE[w], nonces); continue; } + char how[96]; + std::snprintf(how, sizeof(how), "in batch of 2^%d, work-group %d, sub-group width %u", batchLog2, IGNEUM_GROUP, gSubGroupWidth); + overall = compareWarp((const uint64_t*)out.data() + IGNEUM_VEC_BASE[w], IGNEUM_VEC_OUT[w], IGNEUM_VEC_BASE[w], how) && overall; + } + // Fingerprint of the whole batch so different configurations can be compared bit for bit. + std::printf("batch fingerprint (FNV-1a 64 of 2^%d outputs): %016llx\n", batchLog2, (unsigned long long)fnv1a64(out.data(), (size_t)nonces * 8u)); + std::printf("OVERALL: %s\n", overall ? "PASS" : "FAIL"); + return overall ? 0 : 1; +} diff --git a/proto-opencl/emu/emu_opencl.h b/proto-opencl/emu/emu_opencl.h new file mode 100644 index 000000000..cc5edf02e --- /dev/null +++ b/proto-opencl/emu/emu_opencl.h @@ -0,0 +1,49 @@ +// CPU emulation shim: the slice of OpenCL C that the generated kernel.cl uses, as C++. +// kernel.cl includes this file when __OPENCL_VERSION__ is not defined (its own prelude does that), so the exact +// generated text compiles with clang++ or g++ and runs on host threads. emu_main.cpp supplies the runtime below. +// The sub-group width is a runtime setting (32 or 64) so the kernel can be checked as a 64-wide hardware wave +// (AMD GCN/CDNA, RDNA in wave64) would run it: two logical 32-lane units in one wave. Not part of any deliverable +// that runs on a GPU, and it says nothing about AMD hardware; only PASS/FAIL matters, rates are noise. +#pragma once +#include +#include + +typedef unsigned int uint; +typedef unsigned long ulong; +static_assert(sizeof(uint) == 4, "uint must be 32 bits"); +static_assert(sizeof(ulong) == 8, "ulong must be 64 bits (LP64 host: macOS or Linux)"); + +// Address-space qualifiers and the kernel attribute mean nothing on the host. +#define __kernel +#define __global +#define __constant +#define __local +#define __private +#define CLK_LOCAL_MEM_FENCE 1 +#define CLK_GLOBAL_MEM_FENCE 2 +#define IGNEUM_KERNEL_HASH +// Work-group local memory: one arena per work-group instance, shared by its work-items (emu_local_words). +#define IGNEUM_LOCAL_WORDS(name, n) uint* name = emu_local_words((uint)(n)) + +// Work-item functions (thread-local state set by emu_launch). +size_t get_global_id(uint dim); +size_t get_local_id(uint dim); +size_t get_group_id(uint dim); +size_t get_local_size(uint dim); +size_t get_global_size(uint dim); +uint get_sub_group_size(void); +uint get_sub_group_local_id(void); +uint get_sub_group_id(void); +uint get_num_sub_groups(void); + +// Synchronisation and exchange. +void barrier(int flags); // work-group barrier +uint sub_group_shuffle_xor(uint v, uint mask); // cl_khr_subgroup_shuffle semantics over the emulated sub-group +uint intel_sub_group_shuffle_xor(uint v, uint mask); // same semantics +uint sub_group_broadcast(uint v, uint lane); +uint* emu_local_words(uint n); + +// Integer built-ins with OpenCL semantics. +static inline uint mul_hi(uint a, uint b) { return (uint)(((uint64_t)a * (uint64_t)b) >> 32); } +// rotate(v, i): bits shifted left by i modulo the bit width (OpenCL C spec 6.3 and 6.12.3). +static inline uint rotate(uint x, uint n) { n &= 31u; return n == 0u ? x : ((x << n) | (x >> (32u - n))); } diff --git a/proto-opencl/host.c b/proto-opencl/host.c new file mode 100644 index 000000000..6a9b455e7 --- /dev/null +++ b/proto-opencl/host.c @@ -0,0 +1,940 @@ +// igneum-bench-cl: OpenCL host program for Igneum's random-program proof-of-work test harness. +// The portable third path after Apple Metal (proto-metal) and NVIDIA CUDA (proto-cuda): it runs on AMD (Windows +// and Linux), NVIDIA, Intel and, as a correctness check only, on Apple's deprecated OpenCL 1.2 runtime. +// +// TEST HARNESS ONLY. No pool, no network, no wallet, no mining protocol. It fills the dataset on the device, +// checks the device against vectors produced on the Mac (proto-metal), and times the kernel. +// +// C99 plus the OpenCL 1.2 API, nothing else. The kernels are compiled from packs//kernel.cl at runtime. +// The pack's program.h, vectors.h and (memory-hard packs) memhard.h are included at compile time; memhard.h is the +// host reference that fills the cache on one thread and derives dataset words for the self-test. +// +// Build: see README.md (macOS -framework OpenCL, Linux -lOpenCL, Windows cl.exe + OpenCL.lib), or build.sh / build.bat. + +#define _CRT_SECURE_NO_WARNINGS +#define CL_TARGET_OPENCL_VERSION 120 +#define CL_USE_DEPRECATED_OPENCL_1_2_APIS +#if defined(__APPLE__) && !defined(IGNEUM_KHR_HEADERS) +#include +#else +#include +#endif + +#include +#include +#include +#include +#ifdef _WIN32 +#define WIN32_LEAN_AND_MEAN +#include +#else +#include +#include +#endif + +#define IGNEUM_NO_CUDA +#include "program.h" +#include "vectors.h" + +#ifndef IGNEUM_DATASET_MODE +#define IGNEUM_DATASET_MODE 0 +#endif +#if IGNEUM_DATASET_MODE == 1 +#include "memhard.h" +#endif +#ifndef IGNEUM_KERNEL_PATH +#define IGNEUM_KERNEL_PATH "kernel.cl" +#endif + +// Sub-group query constants (cl_khr_subgroups / OpenCL 2.1). Spelled out because OpenCL 1.2 headers lack them. +#define IG_CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE 0x2033 +#define IG_CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE 0x2034 +// Vendor device attributes (cl_amd_device_attribute_query, cl_nv_device_attribute_query). +#define IG_CL_DEVICE_WAVEFRONT_WIDTH_AMD 0x4043 +#define IG_CL_DEVICE_WARP_SIZE_NV 0x4003 + +typedef cl_int (CL_API_CALL *ig_pfn_subgroup_info)(cl_kernel, cl_device_id, cl_uint, size_t, const void*, size_t, void*, size_t*); + +// --------------------------------------------------------------------------------------------- +// Errors and timing + +static const char* clErrName(cl_int e) { + switch (e) { + case CL_SUCCESS: return "CL_SUCCESS"; + case CL_DEVICE_NOT_FOUND: return "CL_DEVICE_NOT_FOUND"; + case CL_DEVICE_NOT_AVAILABLE: return "CL_DEVICE_NOT_AVAILABLE"; + case CL_COMPILER_NOT_AVAILABLE: return "CL_COMPILER_NOT_AVAILABLE"; + case CL_MEM_OBJECT_ALLOCATION_FAILURE: return "CL_MEM_OBJECT_ALLOCATION_FAILURE"; + case CL_OUT_OF_RESOURCES: return "CL_OUT_OF_RESOURCES"; + case CL_OUT_OF_HOST_MEMORY: return "CL_OUT_OF_HOST_MEMORY"; + case CL_PROFILING_INFO_NOT_AVAILABLE: return "CL_PROFILING_INFO_NOT_AVAILABLE"; + case CL_BUILD_PROGRAM_FAILURE: return "CL_BUILD_PROGRAM_FAILURE"; + case CL_INVALID_VALUE: return "CL_INVALID_VALUE"; + case CL_INVALID_DEVICE: return "CL_INVALID_DEVICE"; + case CL_INVALID_CONTEXT: return "CL_INVALID_CONTEXT"; + case CL_INVALID_QUEUE_PROPERTIES: return "CL_INVALID_QUEUE_PROPERTIES"; + case CL_INVALID_COMMAND_QUEUE: return "CL_INVALID_COMMAND_QUEUE"; + case CL_INVALID_MEM_OBJECT: return "CL_INVALID_MEM_OBJECT"; + case CL_INVALID_BUFFER_SIZE: return "CL_INVALID_BUFFER_SIZE"; + case CL_INVALID_BUILD_OPTIONS: return "CL_INVALID_BUILD_OPTIONS"; + case CL_INVALID_PROGRAM: return "CL_INVALID_PROGRAM"; + case CL_INVALID_PROGRAM_EXECUTABLE: return "CL_INVALID_PROGRAM_EXECUTABLE"; + case CL_INVALID_KERNEL_NAME: return "CL_INVALID_KERNEL_NAME"; + case CL_INVALID_KERNEL: return "CL_INVALID_KERNEL"; + case CL_INVALID_ARG_INDEX: return "CL_INVALID_ARG_INDEX"; + case CL_INVALID_ARG_VALUE: return "CL_INVALID_ARG_VALUE"; + case CL_INVALID_ARG_SIZE: return "CL_INVALID_ARG_SIZE"; + case CL_INVALID_KERNEL_ARGS: return "CL_INVALID_KERNEL_ARGS"; + case CL_INVALID_WORK_DIMENSION: return "CL_INVALID_WORK_DIMENSION"; + case CL_INVALID_WORK_GROUP_SIZE: return "CL_INVALID_WORK_GROUP_SIZE"; + case CL_INVALID_WORK_ITEM_SIZE: return "CL_INVALID_WORK_ITEM_SIZE"; + case CL_INVALID_GLOBAL_OFFSET: return "CL_INVALID_GLOBAL_OFFSET"; + case CL_INVALID_EVENT: return "CL_INVALID_EVENT"; + case CL_INVALID_OPERATION: return "CL_INVALID_OPERATION"; + case CL_INVALID_GLOBAL_WORK_SIZE: return "CL_INVALID_GLOBAL_WORK_SIZE"; + case CL_INVALID_PLATFORM: return "CL_INVALID_PLATFORM"; + default: return "(other)"; + } +} + +static void clFail(cl_int e, const char* what, int line) { + fprintf(stderr, "OpenCL error: %s (%d)\n at host.c:%d\n in %s\n", clErrName(e), (int)e, line, what); + exit(2); +} +#define CL_CHECK(call) do { cl_int err_ = (call); if (err_ != CL_SUCCESS) clFail(err_, #call, __LINE__); } while (0) +#define CL_CHECK_ERR(err_, what) do { if ((err_) != CL_SUCCESS) clFail((err_), what, __LINE__); } while (0) + +static double wallMs(void) { +#ifdef _WIN32 + LARGE_INTEGER f, c; + QueryPerformanceFrequency(&f); + QueryPerformanceCounter(&c); + return (double)c.QuadPart * 1000.0 / (double)f.QuadPart; +#else + struct timespec ts; + clock_gettime(CLOCK_MONOTONIC, &ts); + return (double)ts.tv_sec * 1000.0 + (double)ts.tv_nsec / 1e6; +#endif +} + +// Event profiling. A runtime that cannot report timestamps (CL_PROFILING_INFO_NOT_AVAILABLE) gives -1 and the +// harness switches the rate to wall time instead of stopping; the count of such events is reported. +static int gProfilingFailures = 0; +static double eventMs(cl_event e) { + cl_ulong t0 = 0, t1 = 0; + if (clGetEventProfilingInfo(e, CL_PROFILING_COMMAND_START, sizeof(t0), &t0, NULL) != CL_SUCCESS || + clGetEventProfilingInfo(e, CL_PROFILING_COMMAND_END, sizeof(t1), &t1, NULL) != CL_SUCCESS) { ++gProfilingFailures; return -1.0; } + return (double)(t1 - t0) / 1e6; +} + +static double spanMs(cl_event first, cl_event last) { + cl_ulong t0 = 0, t1 = 0; + if (clGetEventProfilingInfo(first, CL_PROFILING_COMMAND_START, sizeof(t0), &t0, NULL) != CL_SUCCESS || + clGetEventProfilingInfo(last, CL_PROFILING_COMMAND_END, sizeof(t1), &t1, NULL) != CL_SUCCESS) { ++gProfilingFailures; return -1.0; } + return (double)(t1 - t0) / 1e6; +} + +static void* loadSym(const char* name) { +#ifdef _WIN32 + HMODULE m = GetModuleHandleA("OpenCL.dll"); + return m ? (void*)GetProcAddress(m, name) : NULL; +#else + return dlsym(RTLD_DEFAULT, name); +#endif +} + +// --------------------------------------------------------------------------------------------- +// Host reference + +static const uint32_t SEEDW[8] = IGNEUM_SEEDW_INIT; + +#if IGNEUM_DATASET_MODE == 0 +// Same closed form as ds_elem in kernel.cl and datasetElem in proto-metal/main.swift. +static uint32_t host_ds_elem(uint32_t i, uint32_t d0, uint32_t d1) { + uint32_t x = i ^ d0; + x *= 0x9E3779B1u; x ^= x >> 15; + x += d1; + x *= 0x85EBCA77u; x ^= x >> 13; + x *= 0xC2B2AE3Du; x ^= x >> 16; + return x; +} +#else +static uint64_t fnv1a64(const void* p, size_t n) { + const uint8_t* b = (const uint8_t*)p; + uint64_t h = 0xcbf29ce484222325ull; + size_t i; + for (i = 0; i < n; ++i) { h ^= b[i]; h *= 0x100000001b3ull; } + return h; +} +static const uint32_t CACHE_WORDS_HOST = 1u << IGNEUM_CACHE_LOG2_WORDS; +static uint32_t* hCache = NULL; +static uint32_t host_ds_word(uint32_t w) { + uint32_t s[16]; + mh_item(hCache, w >> 4u, s); + return s[w & 15u]; +} +#endif + +// --------------------------------------------------------------------------------------------- +// Options + +typedef struct { + int datasetMib; + int batchLog2; + int batches; + int groupWarps; + int sweep; + int device; // flat index into the enumerated list, -1 = first GPU + int exchange; // 0 auto, 1 force local-memory fallback, 2 force sub-group shuffles + int list; + int timeWall; // 1 = rate from wall time, 0 = from device event profiling, -1 = auto (wall on the Apple platform) + const char* kernelPath; + const char* extraOpts; +} Options; + +static int packMib(void) { return (int)(((1ull << IGNEUM_DATASET_LOG2) * 4ull) >> 20); } + +static void usage(void) { + printf( + "igneum-bench-cl [--list] [--device D] [--dataset-mib N] [--sweep] [--batch-log2 24] [--batches 5] [--group-warps 1]\n" + " [--exchange auto|local|subgroup] [--kernel path/to/kernel.cl] [--build-opts \"...\"]\n" + " --list print every OpenCL platform and device, then exit\n" + " --device D device index from the list (default: the first GPU, else device 0)\n" + " --dataset-mib N dataset size in MiB, power of two (default 1024; vectors are only checked at %d MiB)\n" + " --sweep run 4, 64, 256, 512 and 1024 MiB in sequence (same sweep as the Mac and the CUDA harness)\n" + " --batch-log2 B nonces per batch = 2^B (default 24)\n" + " --batches N timed batches after one warm-up batch (default 5)\n" + " --group-warps W 32-lane units per work-group, 1..8 (default 1 = one work-group per unit; sub-group shuffles need 1)\n" + " --exchange M auto (default): sub-group shuffles when the device has them and its sub-group size is 32, else local memory\n" + " local: force the local-memory exchange; subgroup: require sub-group shuffles or fail\n" + " --kernel P path to the pack's kernel.cl (default: the path compiled in, %s)\n" + " --build-opts S extra options appended to clBuildProgram (for example \"-cl-std=CL2.0\")\n" + " --time T event (default): hashes/s from device event profiling, like cudaEvent time; wall: from host wall time.\n" + " Apple's OpenCL runtime reports unusable event timestamps, so wall is the default on the Apple platform.\n", packMib(), IGNEUM_KERNEL_PATH); +} + +static int isPow2(long long v) { return v > 0 && (v & (v - 1)) == 0; } +static int log2u32(uint32_t v) { int n = 0; while (v > 1u) { v >>= 1; ++n; } return n; } + +static Options parseArgs(int argc, char** argv) { + Options o; + int i; + o.datasetMib = 1024; o.batchLog2 = 24; o.batches = 5; o.groupWarps = 1; o.sweep = 0; o.device = -1; + o.exchange = 0; o.list = 0; o.timeWall = -1; o.kernelPath = IGNEUM_KERNEL_PATH; o.extraOpts = ""; + for (i = 1; i < argc; ++i) { + const char* a = argv[i]; + int needs = (strcmp(a, "--dataset-mib") == 0 || strcmp(a, "--batch-log2") == 0 || strcmp(a, "--batches") == 0 || + strcmp(a, "--group-warps") == 0 || strcmp(a, "--device") == 0 || strcmp(a, "--exchange") == 0 || + strcmp(a, "--kernel") == 0 || strcmp(a, "--build-opts") == 0 || strcmp(a, "--time") == 0); + if (needs && i + 1 >= argc) { usage(); exit(2); } + if (strcmp(a, "--dataset-mib") == 0) o.datasetMib = atoi(argv[++i]); + else if (strcmp(a, "--batch-log2") == 0) o.batchLog2 = atoi(argv[++i]); + else if (strcmp(a, "--batches") == 0) o.batches = atoi(argv[++i]); + else if (strcmp(a, "--group-warps") == 0) o.groupWarps = atoi(argv[++i]); + else if (strcmp(a, "--device") == 0) o.device = atoi(argv[++i]); + else if (strcmp(a, "--kernel") == 0) o.kernelPath = argv[++i]; + else if (strcmp(a, "--build-opts") == 0) o.extraOpts = argv[++i]; + else if (strcmp(a, "--time") == 0) { + const char* m = argv[++i]; + if (strcmp(m, "event") == 0) o.timeWall = 0; + else if (strcmp(m, "wall") == 0) o.timeWall = 1; + else { printf("--time must be event or wall\n"); exit(2); } + } + else if (strcmp(a, "--exchange") == 0) { + const char* m = argv[++i]; + if (strcmp(m, "auto") == 0) o.exchange = 0; + else if (strcmp(m, "local") == 0) o.exchange = 1; + else if (strcmp(m, "subgroup") == 0) o.exchange = 2; + else { printf("--exchange must be auto, local or subgroup\n"); exit(2); } + } + else if (strcmp(a, "--sweep") == 0) o.sweep = 1; + else if (strcmp(a, "--list") == 0) o.list = 1; + else if (strcmp(a, "-h") == 0 || strcmp(a, "--help") == 0) { usage(); exit(0); } + else { printf("unknown argument %s\n", a); usage(); exit(2); } + } + if (!isPow2(o.datasetMib) || o.datasetMib < 1 || o.datasetMib > 16384) { printf("--dataset-mib must be a power of two between 1 and 16384\n"); exit(2); } + if (o.batchLog2 < 10 || o.batchLog2 > 28) { printf("--batch-log2 must be between 10 and 28\n"); exit(2); } + if (o.batches < 1) { printf("--batches must be at least 1\n"); exit(2); } + if (o.groupWarps < 1 || o.groupWarps > 8) { printf("--group-warps must be between 1 and 8\n"); exit(2); } + return o; +} + +// --------------------------------------------------------------------------------------------- +// Devices + +typedef struct { + cl_platform_id platform; + cl_device_id device; + char platformName[256], platformVersion[256]; + char name[256], vendor[256], version[256], driver[256], cVersion[256]; + char* extensions; + cl_device_type type; + cl_uint computeUnits, clockMHz; + cl_ulong globalMem, maxAlloc, localMem; + size_t maxWorkGroup; + int cMajor, cMinor; // OpenCL C version + int dMajor, dMinor; // device (platform profile) version + cl_uint amdWavefront, nvWarp; // 0 if not reported +} DeviceInfo; + +static void devStr(cl_device_id d, cl_device_info what, char* out, size_t n) { + out[0] = 0; + clGetDeviceInfo(d, what, n - 1, out, NULL); + out[n - 1] = 0; +} + +static const char* typeName(cl_device_type t) { + if (t & CL_DEVICE_TYPE_GPU) return "GPU"; + if (t & CL_DEVICE_TYPE_CPU) return "CPU"; + if (t & CL_DEVICE_TYPE_ACCELERATOR) return "accelerator"; + return "other"; +} + +static int enumerateDevices(DeviceInfo** outList) { + cl_uint np = 0, p; + cl_platform_id plats[16]; + DeviceInfo* list = NULL; + int n = 0; + cl_int e = clGetPlatformIDs(16, plats, &np); + if (e != CL_SUCCESS || np == 0) { *outList = NULL; return 0; } + for (p = 0; p < np; ++p) { + cl_uint nd = 0, d; + cl_device_id devs[32]; + char pname[256] = {0}, pver[256] = {0}; + clGetPlatformInfo(plats[p], CL_PLATFORM_NAME, sizeof(pname) - 1, pname, NULL); + clGetPlatformInfo(plats[p], CL_PLATFORM_VERSION, sizeof(pver) - 1, pver, NULL); + if (clGetDeviceIDs(plats[p], CL_DEVICE_TYPE_ALL, 32, devs, &nd) != CL_SUCCESS) continue; + for (d = 0; d < nd; ++d) { + DeviceInfo di; + size_t extLen = 0; + memset(&di, 0, sizeof(di)); + di.platform = plats[p]; di.device = devs[d]; + strncpy(di.platformName, pname, 255); strncpy(di.platformVersion, pver, 255); + devStr(devs[d], CL_DEVICE_NAME, di.name, sizeof(di.name)); + devStr(devs[d], CL_DEVICE_VENDOR, di.vendor, sizeof(di.vendor)); + devStr(devs[d], CL_DEVICE_VERSION, di.version, sizeof(di.version)); + devStr(devs[d], CL_DRIVER_VERSION, di.driver, sizeof(di.driver)); + devStr(devs[d], CL_DEVICE_OPENCL_C_VERSION, di.cVersion, sizeof(di.cVersion)); + clGetDeviceInfo(devs[d], CL_DEVICE_EXTENSIONS, 0, NULL, &extLen); + di.extensions = (char*)calloc(extLen + 1, 1); + if (extLen) clGetDeviceInfo(devs[d], CL_DEVICE_EXTENSIONS, extLen, di.extensions, NULL); + clGetDeviceInfo(devs[d], CL_DEVICE_TYPE, sizeof(di.type), &di.type, NULL); + clGetDeviceInfo(devs[d], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(di.computeUnits), &di.computeUnits, NULL); + clGetDeviceInfo(devs[d], CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof(di.clockMHz), &di.clockMHz, NULL); + clGetDeviceInfo(devs[d], CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(di.globalMem), &di.globalMem, NULL); + clGetDeviceInfo(devs[d], CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(di.maxAlloc), &di.maxAlloc, NULL); + clGetDeviceInfo(devs[d], CL_DEVICE_LOCAL_MEM_SIZE, sizeof(di.localMem), &di.localMem, NULL); + clGetDeviceInfo(devs[d], CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(di.maxWorkGroup), &di.maxWorkGroup, NULL); + if (sscanf(di.cVersion, "OpenCL C %d.%d", &di.cMajor, &di.cMinor) != 2) { di.cMajor = 1; di.cMinor = 2; } + if (sscanf(di.version, "OpenCL %d.%d", &di.dMajor, &di.dMinor) != 2) { di.dMajor = 1; di.dMinor = 2; } + if (strstr(di.extensions, "cl_amd_device_attribute_query")) + clGetDeviceInfo(devs[d], IG_CL_DEVICE_WAVEFRONT_WIDTH_AMD, sizeof(di.amdWavefront), &di.amdWavefront, NULL); + if (strstr(di.extensions, "cl_nv_device_attribute_query")) + clGetDeviceInfo(devs[d], IG_CL_DEVICE_WARP_SIZE_NV, sizeof(di.nvWarp), &di.nvWarp, NULL); + list = (DeviceInfo*)realloc(list, sizeof(DeviceInfo) * (size_t)(n + 1)); + list[n++] = di; + } + } + *outList = list; + return n; +} + +static void printDevice(int idx, const DeviceInfo* d, int chosen) { + const char* subExt = strstr(d->extensions, "cl_khr_subgroup_shuffle") ? "cl_khr_subgroup_shuffle" : + strstr(d->extensions, "cl_intel_subgroups") ? "cl_intel_subgroups" : + strstr(d->extensions, "cl_khr_subgroups") ? "cl_khr_subgroups (no shuffle extension)" : "none"; + printf("%s[%d] %s | %s (%s)\n", chosen ? "*" : " ", idx, d->name, d->platformName, d->platformVersion); + printf(" %s, vendor %s, driver %s, %s, %u compute units, %u MHz\n", typeName(d->type), d->vendor, d->driver, d->cVersion, d->computeUnits, d->clockMHz); + printf(" global %llu MiB, max alloc %llu MiB, local %llu KiB, max work-group %llu, sub-group extension: %s", + (unsigned long long)(d->globalMem >> 20), (unsigned long long)(d->maxAlloc >> 20), (unsigned long long)(d->localMem >> 10), + (unsigned long long)d->maxWorkGroup, subExt); + if (d->amdWavefront) printf(", AMD wavefront width %u", d->amdWavefront); + if (d->nvWarp) printf(", NVIDIA warp size %u", d->nvWarp); + printf("\n"); +} + +// --------------------------------------------------------------------------------------------- +// Program build + +static char* readFile(const char* path, size_t* len) { + FILE* f = fopen(path, "rb"); + char* buf; + long n; + if (!f) return NULL; + fseek(f, 0, SEEK_END); n = ftell(f); fseek(f, 0, SEEK_SET); + if (n < 0) { fclose(f); return NULL; } + buf = (char*)malloc((size_t)n + 1); + if (fread(buf, 1, (size_t)n, f) != (size_t)n) { fclose(f); free(buf); return NULL; } + buf[n] = 0; + fclose(f); + *len = (size_t)n; + return buf; +} + +typedef struct { + cl_context ctx; + cl_command_queue q; + cl_program prog; + cl_kernel kHash, kCacheFill, kBuild, kFill; + int exchange; // 0 local memory, 1 khr sub-group shuffle, 2 intel + size_t subGroupSize; // as queried for a 32-item work-group, 0 if not queried + char exchangeNote[512]; + char buildOptions[512]; +} Device; + +static const char* exchangeName(int m) { return m == 1 ? "sub_group_shuffle_xor (cl_khr_subgroup_shuffle)" : m == 2 ? "intel_sub_group_shuffle_xor (cl_intel_subgroups)" : "local-memory exchange with barrier"; } + +// Returns 0 on success, 1 on build failure (log printed). +static int buildProgram(Device* dv, const DeviceInfo* di, const char* src, size_t srcLen, int exchangeMode, int groupSize, const char* extra) { + cl_int err = 0; + const char* std; + // The sub-group built-ins need OpenCL C 2.0 or 3.0. OpenCL 3.0 devices may report "OpenCL C 1.2" as the default + // CL_DEVICE_OPENCL_C_VERSION while supporting 3.0 (the 3.0 API lists all versions; the 1.2 API cannot ask), so + // the device version counts as well. The local-memory variant is always built as OpenCL C 1.2, the same text everywhere. + int major = di->cMajor > di->dMajor ? di->cMajor : di->dMajor; + if (exchangeMode == 0) std = "-cl-std=CL1.2"; + else if (major >= 3) std = "-cl-std=CL3.0"; + else if (major >= 2) std = "-cl-std=CL2.0"; + else std = "-cl-std=CL1.2"; + snprintf(dv->buildOptions, sizeof(dv->buildOptions), "%s -D IGNEUM_GROUP=%d -D IGNEUM_EXCHANGE=%d %s", std, groupSize, exchangeMode, extra); + dv->prog = clCreateProgramWithSource(dv->ctx, 1, &src, &srcLen, &err); + CL_CHECK_ERR(err, "clCreateProgramWithSource"); + err = clBuildProgram(dv->prog, 1, &di->device, dv->buildOptions, NULL, NULL); + if (err != CL_SUCCESS) { + size_t logLen = 0; + char* log; + clGetProgramBuildInfo(dv->prog, di->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &logLen); + log = (char*)calloc(logLen + 1, 1); + if (logLen) clGetProgramBuildInfo(dv->prog, di->device, CL_PROGRAM_BUILD_LOG, logLen, log, NULL); + printf("build FAILED (%s) with options \"%s\"\n--- build log ---\n%s\n--- end of build log ---\n", clErrName(err), dv->buildOptions, log); + free(log); + clReleaseProgram(dv->prog); dv->prog = NULL; + return 1; + } + dv->kHash = clCreateKernel(dv->prog, "igneum_hash", &err); CL_CHECK_ERR(err, "clCreateKernel igneum_hash"); +#if IGNEUM_DATASET_MODE == 1 + dv->kCacheFill = clCreateKernel(dv->prog, "igneum_cache_fill", &err); CL_CHECK_ERR(err, "clCreateKernel igneum_cache_fill"); + dv->kBuild = clCreateKernel(dv->prog, "igneum_build", &err); CL_CHECK_ERR(err, "clCreateKernel igneum_build"); +#else + dv->kFill = clCreateKernel(dv->prog, "igneum_fill", &err); CL_CHECK_ERR(err, "clCreateKernel igneum_fill"); +#endif + return 0; +} + +static void releaseProgram(Device* dv) { + if (dv->kHash) clReleaseKernel(dv->kHash); + if (dv->kCacheFill) clReleaseKernel(dv->kCacheFill); + if (dv->kBuild) clReleaseKernel(dv->kBuild); + if (dv->kFill) clReleaseKernel(dv->kFill); + if (dv->prog) clReleaseProgram(dv->prog); + dv->kHash = dv->kCacheFill = dv->kBuild = dv->kFill = NULL; dv->prog = NULL; +} + +// Sub-group size of igneum_hash for a work-group of `local` items. 0 if the query is unavailable (reason in *why). +static size_t querySubGroupSize(const Device* dv, const DeviceInfo* di, size_t local, char* why, size_t whyLen) { + ig_pfn_subgroup_info fn = (ig_pfn_subgroup_info)clGetExtensionFunctionAddressForPlatform(di->platform, "clGetKernelSubGroupInfoKHR"); + const char* via = "clGetKernelSubGroupInfoKHR"; + size_t sg = 0; + cl_int e; + if (!fn) { fn = (ig_pfn_subgroup_info)loadSym("clGetKernelSubGroupInfo"); via = "clGetKernelSubGroupInfo (OpenCL 2.1 core, through the loader)"; } + if (!fn) { snprintf(why, whyLen, "the sub-group size could not be queried (neither clGetKernelSubGroupInfoKHR nor clGetKernelSubGroupInfo is available)"); return 0; } + e = fn(dv->kHash, di->device, IG_CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE, sizeof(local), &local, sizeof(sg), &sg, NULL); + if (e != CL_SUCCESS) { snprintf(why, whyLen, "the sub-group size query failed: %s returned %s (%d)", via, clErrName(e), (int)e); return 0; } + snprintf(why, whyLen, "queried through %s", via); + return sg; +} + +// Decide the exchange implementation and build. See WAVEFRONT.md for the rule. +static void setupProgram(Device* dv, const DeviceInfo* di, const Options* o, const char* src, size_t srcLen) { + int groupSize = 32 * o->groupWarps; + int want = 0; + const char* why = ""; + if (strstr(di->extensions, "cl_khr_subgroup_shuffle")) want = 1; + else if (strstr(di->extensions, "cl_intel_subgroups")) want = 2; + else why = "device lists no sub-group shuffle extension"; + if (o->exchange == 1) { want = 0; why = "forced by --exchange local"; } + if (want != 0 && o->groupWarps != 1) { + if (o->exchange == 2) { printf("FAIL: --exchange subgroup needs --group-warps 1 (one work-group = one 32-lane unit)\n"); exit(2); } + want = 0; why = "--group-warps is not 1, so a work-group is not one 32-lane unit"; + } + if (want == 0 && o->exchange == 2) { printf("FAIL: --exchange subgroup requested but %s\n", why); exit(2); } + + if (want != 0) { + if (buildProgram(dv, di, src, srcLen, want, groupSize, o->extraOpts) != 0) { + if (o->exchange == 2) { printf("FAIL: the sub-group variant did not compile\n"); exit(2); } + want = 0; why = "the sub-group variant did not compile (log above)"; + } else { + static char qwhy[256]; + size_t sg = querySubGroupSize(dv, di, 32, qwhy, sizeof(qwhy)); + if (sg == 0) { + // Second choice: a probe kernel from the same build reports get_sub_group_size() for a 32-item work-group. + // Weaker than the per-kernel query (a compiler may pick the wave width per kernel), and the note says so. + cl_int perr = 0; + cl_kernel kp = clCreateKernel(dv->prog, "igneum_probe_subgroup", &perr); + if (perr == CL_SUCCESS) { + cl_uint probe[2] = { 0u, 0u }; + cl_mem pb = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, sizeof(probe), NULL, &perr); + size_t g = 32, l = 32; + if (perr == CL_SUCCESS && clSetKernelArg(kp, 0, sizeof(cl_mem), &pb) == CL_SUCCESS && + clEnqueueNDRangeKernel(dv->q, kp, 1, NULL, &g, &l, 0, NULL, NULL) == CL_SUCCESS && + clEnqueueReadBuffer(dv->q, pb, CL_TRUE, 0, sizeof(probe), probe, 0, NULL, NULL) == CL_SUCCESS && probe[0] != 0u) { + static char pwhy[400]; + snprintf(pwhy, sizeof(pwhy), "%s; probe kernel reports get_sub_group_size() %u and %u sub-group(s) for a 32-item work-group (weaker than the per-kernel query)", qwhy, probe[0], probe[1]); + sg = probe[0]; + strncpy(qwhy, pwhy, sizeof(qwhy) - 1); qwhy[sizeof(qwhy) - 1] = 0; + } + if (pb) clReleaseMemObject(pb); + clReleaseKernel(kp); + } + } + dv->subGroupSize = sg; + if (sg == 32) { + dv->exchange = want; + snprintf(dv->exchangeNote, sizeof(dv->exchangeNote), "%s, sub-group size %llu for a 32-item work-group (%s)", exchangeName(want), (unsigned long long)sg, qwhy); + return; + } + releaseProgram(dv); + if (sg == 0) why = qwhy; + else why = "the sub-group size for a 32-item work-group is not 32"; + if (o->exchange == 2) { printf("FAIL: --exchange subgroup requested but %s (sub-group size %llu)\n", why, (unsigned long long)sg); exit(2); } + } + } + if (buildProgram(dv, di, src, srcLen, 0, groupSize, o->extraOpts) != 0) { printf("FAIL: kernel.cl did not compile\n"); exit(2); } + dv->exchange = 0; + if (dv->subGroupSize) snprintf(dv->exchangeNote, sizeof(dv->exchangeNote), "%s (%s; queried sub-group size %llu)", exchangeName(0), why, (unsigned long long)dv->subGroupSize); + else snprintf(dv->exchangeNote, sizeof(dv->exchangeNote), "%s (%s)", exchangeName(0), why); +} + +// --------------------------------------------------------------------------------------------- +// Launch helpers + +static size_t kernelMaxLocal(const Device* dv, cl_kernel k, const DeviceInfo* di, size_t want) { + size_t wg = 0; + (void)dv; + if (clGetKernelWorkGroupInfo(k, di->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(wg), &wg, NULL) != CL_SUCCESS || wg == 0) wg = want; + if (wg > di->maxWorkGroup) wg = di->maxWorkGroup; + return wg < want ? wg : want; +} + +static cl_event launch1D(const Device* dv, cl_kernel k, size_t global, size_t local) { + cl_event ev = NULL; + size_t g = ((global + local - 1) / local) * local; + CL_CHECK(clEnqueueNDRangeKernel(dv->q, k, 1, NULL, &g, &local, 0, NULL, &ev)); + return ev; +} + +static cl_event launchHash(const Device* dv, cl_mem ds, cl_mem out, cl_uint baseNonce, cl_uint mask, size_t nonces, size_t groupSize) { + CL_CHECK(clSetKernelArg(dv->kHash, 0, sizeof(cl_mem), &ds)); + CL_CHECK(clSetKernelArg(dv->kHash, 1, sizeof(cl_mem), &out)); + CL_CHECK(clSetKernelArg(dv->kHash, 2, sizeof(cl_uint), &baseNonce)); + CL_CHECK(clSetKernelArg(dv->kHash, 3, sizeof(cl_uint), &mask)); + return launch1D(dv, dv->kHash, nonces, groupSize); +} + +static void readWords(const Device* dv, cl_mem buf, size_t wordIndex, size_t nWords, uint32_t* dst) { + CL_CHECK(clEnqueueReadBuffer(dv->q, buf, CL_TRUE, wordIndex * 4u, nWords * 4u, dst, 0, NULL, NULL)); +} + +static int compareWarp(const uint64_t* got, const uint64_t* want, uint32_t base, const char* how) { + int bad = 0, first = -1, l; + for (l = 0; l < 32; ++l) if (got[l] != want[l]) { if (first < 0) first = l; ++bad; } + if (bad == 0) printf("verify warp base %u (nonces %u..%u) %s: PASS\n", base, base, base + 31u, how); + else printf("verify warp base %u (nonces %u..%u) %s: FAIL %d of 32 lanes differ, first lane %d: device=%016llx expected=%016llx\n", + base, base, base + 31u, how, bad, first, (unsigned long long)got[first], (unsigned long long)want[first]); + return bad == 0; +} + +// --------------------------------------------------------------------------------------------- +// Cache (memory-hard packs) + +#if IGNEUM_DATASET_MODE == 1 +static cl_mem gCache = NULL; +static double gCacheFillFirstMs = 0, gCacheFillSecondMs = 0, gCacheHostMs = 0; +static int gCachePass = 0; + +static int setupCache(Device* dv, const DeviceInfo* di) { + size_t bytes = (size_t)CACHE_WORDS_HOST * 4u; + cl_int err = 0; + cl_uint nSeg = IGNEUM_CACHE_SEGMENTS; + size_t local; + int pass; + uint32_t* dev; + uint32_t seg; + double h0; + int same, fnvOk, headOk, lastOk; + uint64_t fnv; + gCache = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, bytes, NULL, &err); + CL_CHECK_ERR(err, "clCreateBuffer cache"); + local = kernelMaxLocal(dv, dv->kCacheFill, di, 256); + for (pass = 0; pass < 2; ++pass) { + cl_event ev; + CL_CHECK(clSetKernelArg(dv->kCacheFill, 0, sizeof(cl_mem), &gCache)); + CL_CHECK(clSetKernelArg(dv->kCacheFill, 1, sizeof(cl_uint), &nSeg)); + ev = launch1D(dv, dv->kCacheFill, nSeg, local); + CL_CHECK(clFinish(dv->q)); + if (pass == 0) gCacheFillFirstMs = eventMs(ev); else gCacheFillSecondMs = eventMs(ev); + clReleaseEvent(ev); + } + printf("cache fill (device): %.2f ms first, %.2f ms second (%u chains x %u ChaCha blocks, %u MiB, work-group %llu)\n", + gCacheFillFirstMs, gCacheFillSecondMs, (unsigned)IGNEUM_CACHE_SEGMENTS, 1u << IGNEUM_CACHE_SEGMENT_LOG2_LINES, + (unsigned)(bytes >> 20), (unsigned long long)local); + + hCache = (uint32_t*)malloc(bytes); + h0 = wallMs(); + for (seg = 0; seg < IGNEUM_CACHE_SEGMENTS; ++seg) mh_cache_segment(hCache, seg); + gCacheHostMs = wallMs() - h0; + printf("cache fill (host, one thread): %.1f ms\n", gCacheHostMs); + + dev = (uint32_t*)malloc(bytes); + readWords(dv, gCache, 0, CACHE_WORDS_HOST, dev); + same = memcmp(dev, hCache, bytes) == 0; + fnv = fnv1a64(hCache, bytes); + fnvOk = (fnv == IGNEUM_CACHE_FNV64); + headOk = memcmp(hCache, IGNEUM_CACHE_HEAD, 64) == 0; + lastOk = memcmp(hCache + CACHE_WORDS_HOST - 16u, IGNEUM_CACHE_LAST, 64) == 0; + if (!same) { + uint32_t i; + for (i = 0; i < CACHE_WORDS_HOST; ++i) if (dev[i] != hCache[i]) { + printf(" cache[%u]: device 0x%08x host 0x%08x (first difference)\n", i, dev[i], hCache[i]); break; + } + } + free(dev); + gCachePass = same && fnvOk && headOk && lastOk; + printf("cache check: %s (device == host all %u words %s, host FNV-1a 64 %016llx vs Mac %016llx %s, head 16 vs Mac %s, last line vs Mac %s)\n", + gCachePass ? "PASS" : "FAIL", CACHE_WORDS_HOST, same ? "PASS" : "FAIL", + (unsigned long long)fnv, (unsigned long long)IGNEUM_CACHE_FNV64, fnvOk ? "PASS" : "FAIL", + headOk ? "PASS" : "FAIL", lastOk ? "PASS" : "FAIL"); + return gCachePass; +} +#endif + +// --------------------------------------------------------------------------------------------- +// One dataset size: fill or build, self-test, vectors, bench + +typedef struct { + int mib; + uint32_t words; + double fillFirstMs, fillSecondMs; + int dsPass; + int vecChecked, vecPass; + double gpuMs, wallMsTimed; + double hashesPerSec, gbps; +} SizeResult; + +static SizeResult runSize(Device* dv, const DeviceInfo* di, const Options* o, int mib, cl_mem dOut, uint32_t nonces) { + SizeResult r; + uint64_t bytes = (uint64_t)mib << 20; + uint32_t mask; + int atPackSize; + cl_int err = 0; + cl_mem dDs; + int pass; + size_t groupSize = 32u * (size_t)o->groupWarps; + uint64_t got[32]; + double w0, w1, t0, t1, total; + cl_event* evs; + int b; + memset(&r, 0, sizeof(r)); + r.mib = mib; + r.words = (uint32_t)(bytes / 4ull); + mask = r.words - 1u; + atPackSize = (r.words == (1u << IGNEUM_DATASET_LOG2)); + printf("\n=== dataset %d MiB (2^%d words, mask 0x%08x)%s ===\n", mib, log2u32(r.words), mask, + atPackSize ? "" : " [not the pack size: vectors skipped, dataset head and random points still checked]"); + if ((uint64_t)di->maxAlloc < bytes) { + printf("FAIL: CL_DEVICE_MAX_MEM_ALLOC_SIZE is %llu MiB, the dataset needs %d MiB in one buffer\n", (unsigned long long)(di->maxAlloc >> 20), mib); + exit(2); + } + dDs = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, (size_t)bytes, NULL, &err); + CL_CHECK_ERR(err, "clCreateBuffer dataset"); + + // Fill (or build) twice: the Mac showed a first-touch cost on the first fill of a process. + for (pass = 0; pass < 2; ++pass) { + cl_event ev; +#if IGNEUM_DATASET_MODE == 1 + cl_uint nItems = r.words / 16u; + size_t local = kernelMaxLocal(dv, dv->kBuild, di, 256); + CL_CHECK(clSetKernelArg(dv->kBuild, 0, sizeof(cl_mem), &dDs)); + CL_CHECK(clSetKernelArg(dv->kBuild, 1, sizeof(cl_mem), &gCache)); + CL_CHECK(clSetKernelArg(dv->kBuild, 2, sizeof(cl_uint), &nItems)); + ev = launch1D(dv, dv->kBuild, nItems, local); +#else + cl_uint n = r.words, d0 = IGNEUM_DAY0, d1 = IGNEUM_DAY1; + size_t local = kernelMaxLocal(dv, dv->kFill, di, 256); + CL_CHECK(clSetKernelArg(dv->kFill, 0, sizeof(cl_mem), &dDs)); + CL_CHECK(clSetKernelArg(dv->kFill, 1, sizeof(cl_uint), &n)); + CL_CHECK(clSetKernelArg(dv->kFill, 2, sizeof(cl_uint), &d0)); + CL_CHECK(clSetKernelArg(dv->kFill, 3, sizeof(cl_uint), &d1)); + ev = launch1D(dv, dv->kFill, n, local); +#endif + CL_CHECK(clFinish(dv->q)); + if (pass == 0) r.fillFirstMs = eventMs(ev); else r.fillSecondMs = eventMs(ev); + clReleaseEvent(ev); + } +#if IGNEUM_DATASET_MODE == 1 + printf("dataset build (memory-hard, from the cache): %.2f ms first, %.2f ms second -> %.1f M items/s, %.2f G cache-line reads/s (second, device time)\n", + r.fillFirstMs, r.fillSecondMs, (double)(r.words / 16u) / 1e6 / (r.fillSecondMs / 1000.0), + (double)(r.words / 16u) * (double)IGNEUM_ITEM_ROUNDS / 1e9 / (r.fillSecondMs / 1000.0)); +#else + printf("dataset fill: %.2f ms first, %.2f ms second -> %.0f GB/s write (second, device time)\n", + r.fillFirstMs, r.fillSecondMs, (double)bytes / 1e9 / (r.fillSecondMs / 1000.0)); +#endif + + // Dataset self-test: head 16 (any size), element [MASK] (pack size only), 64 pseudo-random points vs the host + // formula (closed form) or the host derivation from the host cache (memory-hard), and the Mac's 64 sampled words. + { + uint32_t head[16]; + int badHead = 0, i, k, badRnd = 0, badSample = 0, nSample = 0, lastOk = 1; + const char* lastText = "skipped"; + uint64_t s = 0x9E3779B97F4A7C15ull ^ (uint64_t)r.words; + readWords(dv, dDs, 0, 16, head); + for (i = 0; i < 16; ++i) if (head[i] != IGNEUM_DS_HEAD[i]) { + if (badHead == 0) printf(" dataset[%d] = 0x%08x, expected 0x%08x\n", i, head[i], IGNEUM_DS_HEAD[i]); + ++badHead; + } + if (atPackSize) { + uint32_t last = 0; + readWords(dv, dDs, IGNEUM_DS_LAST_INDEX, 1, &last); + lastOk = (last == IGNEUM_DS_LAST); + lastText = lastOk ? "PASS" : "FAIL"; + if (!lastOk) printf(" dataset[%u] = 0x%08x, expected 0x%08x\n", IGNEUM_DS_LAST_INDEX, last, IGNEUM_DS_LAST); + } + for (k = 0; k < 64; ++k) { + uint64_t z; + uint32_t idx, v = 0, want; + s += 0x9E3779B97F4A7C15ull; + z = s; + z = (z ^ (z >> 30)) * 0xBF58476D1CE4E5B9ull; + z = (z ^ (z >> 27)) * 0x94D049BB133111EBull; + z ^= z >> 31; + idx = (uint32_t)z & mask; + readWords(dv, dDs, idx, 1, &v); +#if IGNEUM_DATASET_MODE == 1 + want = host_ds_word(idx); +#else + want = host_ds_elem(idx, IGNEUM_DAY0, IGNEUM_DAY1); +#endif + if (v != want) { + if (badRnd == 0) printf(" dataset[%u] = 0x%08x, host %s 0x%08x\n", idx, v, IGNEUM_DATASET_MODE == 1 ? "derivation" : "formula", want); + ++badRnd; + } + } +#ifdef IGNEUM_DS_SAMPLES + for (k = 0; k < IGNEUM_DS_SAMPLES; ++k) { + uint32_t v = 0; + if (IGNEUM_DS_SAMPLE_INDEX[k] > mask) continue; + ++nSample; + readWords(dv, dDs, IGNEUM_DS_SAMPLE_INDEX[k], 1, &v); + if (v != IGNEUM_DS_SAMPLE_VALUE[k]) { + if (badSample == 0) printf(" dataset[%u] = 0x%08x, Mac 0x%08x\n", IGNEUM_DS_SAMPLE_INDEX[k], v, IGNEUM_DS_SAMPLE_VALUE[k]); + ++badSample; + } + } +#endif + r.dsPass = (badHead == 0 && lastOk && badRnd == 0 && badSample == 0); + printf("dataset self-test: %s (head 16 vs Mac %s, element [MASK] vs Mac %s, 64 random points vs host %s %s, %d Mac samples %s)\n", + r.dsPass ? "PASS" : "FAIL", badHead == 0 ? "PASS" : "FAIL", lastText, + IGNEUM_DATASET_MODE == 1 ? "derivation" : "formula", badRnd == 0 ? "PASS" : "FAIL", + nSample, nSample == 0 ? "none in pack" : (badSample == 0 ? "PASS" : "FAIL")); + } + + // Vectors, standalone: one 32-item work-group per base nonce, exactly like the Mac cross-check. + if (atPackSize) { + int w; + r.vecChecked = 1; r.vecPass = 1; + for (w = 0; w < IGNEUM_VEC_WARPS; ++w) { + cl_event ev; + if (o->groupWarps != 1) { + // reqd_work_group_size pins the hash kernel to 32 x group-warps items; run one full group and read its first 32. + ev = launchHash(dv, dDs, dOut, IGNEUM_VEC_BASE[w], mask, groupSize, groupSize); + } else { + ev = launchHash(dv, dDs, dOut, IGNEUM_VEC_BASE[w], mask, 32u, 32u); + } + CL_CHECK(clFinish(dv->q)); + clReleaseEvent(ev); + CL_CHECK(clEnqueueReadBuffer(dv->q, dOut, CL_TRUE, 0, sizeof(got), got, 0, NULL, NULL)); + r.vecPass = compareWarp(got, IGNEUM_VEC_OUT[w], IGNEUM_VEC_BASE[w], o->groupWarps == 1 ? "standalone, 1 unit/work-group" : "standalone, first unit of one work-group") && r.vecPass; + } + } else { + printf("vectors: skipped (the pack's vectors are for %d MiB)\n", packMib()); + } + + // Warm-up batch at base nonce 0. With the default 2^24 nonces it contains all three vector warps, + // so the bench configuration itself (work-group = 32 x group-warps) is also checked bit for bit. + w0 = wallMs(); + { + cl_event ev = launchHash(dv, dDs, dOut, 0u, mask, nonces, groupSize); + CL_CHECK(clFinish(dv->q)); + clReleaseEvent(ev); + } + w1 = wallMs(); + printf("warm-up batch: %u hashes in %.2f ms wall\n", nonces, w1 - w0); + if (atPackSize) { + // Fingerprint of every output in the batch, so two runs (two devices, two exchange paths, the CPU emulator at the + // same --batch-log2) can be compared for all nonces, not only the vector warps. + uint64_t* all = (uint64_t*)malloc((size_t)nonces * sizeof(uint64_t)); + uint64_t fp = 0xcbf29ce484222325ull; + size_t i; + CL_CHECK(clEnqueueReadBuffer(dv->q, dOut, CL_TRUE, 0, (size_t)nonces * sizeof(uint64_t), all, 0, NULL, NULL)); + for (i = 0; i < (size_t)nonces * 8u; ++i) { fp ^= ((const uint8_t*)all)[i]; fp *= 0x100000001b3ull; } + free(all); + printf("batch fingerprint (FNV-1a 64 of 2^%d outputs at base nonce 0): %016llx\n", o->batchLog2, (unsigned long long)fp); + } + if (atPackSize) { + char how[64]; + int w; + snprintf(how, sizeof(how), "in batch, %d unit(s)/work-group", o->groupWarps); + for (w = 0; w < IGNEUM_VEC_WARPS; ++w) { + if ((uint64_t)IGNEUM_VEC_BASE[w] + 32ull > (uint64_t)nonces) { + printf("verify warp base %u in batch: skipped (batch has %u nonces)\n", IGNEUM_VEC_BASE[w], nonces); + continue; + } + CL_CHECK(clEnqueueReadBuffer(dv->q, dOut, CL_TRUE, (size_t)IGNEUM_VEC_BASE[w] * 8u, sizeof(got), got, 0, NULL, NULL)); + r.vecPass = compareWarp(got, IGNEUM_VEC_OUT[w], IGNEUM_VEC_BASE[w], how) && r.vecPass; + } + } + + // Timed batches. Base nonces (b * nonces) mod 2^32, as in the Metal and CUDA runs. Device time is the span from + // the first batch's start to the last batch's end from event profiling, like cudaEvent elapsed time. + evs = (cl_event*)calloc((size_t)o->batches, sizeof(cl_event)); + t0 = wallMs(); + for (b = 1; b <= o->batches; ++b) { + uint32_t base = (uint32_t)((uint64_t)b * (uint64_t)nonces); + evs[b - 1] = launchHash(dv, dDs, dOut, base, mask, nonces, groupSize); + } + CL_CHECK(clFinish(dv->q)); + t1 = wallMs(); + total = (double)nonces * (double)o->batches; + r.gpuMs = spanMs(evs[0], evs[o->batches - 1]); + for (b = 0; b < o->batches; ++b) clReleaseEvent(evs[b]); + free(evs); + r.wallMsTimed = t1 - t0; + if (r.gpuMs < 0 && !o->timeWall) { printf("NOTE: event profiling unavailable on this runtime; the rate uses wall time\n"); ((Options*)o)->timeWall = 1; } + r.hashesPerSec = total / ((o->timeWall ? r.wallMsTimed : r.gpuMs) / 1000.0); + r.gbps = r.hashesPerSec * (double)IGNEUM_LOADS_PER_HASH * 4.0 / 1e9; + printf("timed: %d batches x %u hashes = %.0f hashes (rate below from %s time)\n", o->batches, nonces, total, o->timeWall ? "wall" : "device event"); + printf(" device %.2f ms -> %.3f Mhash/s%s\n", r.gpuMs, total / (r.gpuMs / 1000.0) / 1e6, o->timeWall ? " (event profiling, not used for the rate on this platform)" : ""); + printf(" wall %.2f ms -> %.3f Mhash/s\n", r.wallMsTimed, total / (r.wallMsTimed / 1000.0) / 1e6); + printf(" rate %.3f Mhash/s (%.0f hashes/s), %.2f GB/s useful (loads x 4 B)\n", r.hashesPerSec / 1e6, r.hashesPerSec, r.gbps); + + CL_CHECK(clReleaseMemObject(dDs)); + return r; +} + +// --------------------------------------------------------------------------------------------- +// Main + +int main(int argc, char** argv) { + Options o = parseArgs(argc, argv); + DeviceInfo* devs = NULL; + int nDev, i, chosen = -1; + DeviceInfo* di; + Device dv; + cl_int err = 0; + size_t srcLen = 0; + char* src; + uint32_t nonces; + cl_mem dOut; + int sizes[5], nSizes = 0; + SizeResult results[5]; + int cachePass = 1, overall, anyVec = 0; + + printf("igneum-bench-cl pack \"%s\" (test harness: no pool, no network, no wallet)\n", IGNEUM_SEED_STRING); + nDev = enumerateDevices(&devs); + if (nDev == 0) { printf("FAIL: no OpenCL platform or device found (is an OpenCL driver / ICD installed?)\n"); return 2; } + if (o.device >= 0) { + if (o.device >= nDev) { printf("FAIL: --device %d out of range (%d devices)\n", o.device, nDev); return 2; } + chosen = o.device; + } else { + for (i = 0; i < nDev; ++i) if (devs[i].type & CL_DEVICE_TYPE_GPU) { chosen = i; break; } + if (chosen < 0) chosen = 0; + } + printf("OpenCL devices (%d):\n", nDev); + for (i = 0; i < nDev; ++i) printDevice(i, &devs[i], i == chosen && !o.list); + if (o.list) return 0; + di = &devs[chosen]; + printf("using device [%d] %s\n", chosen, di->name); + if (o.timeWall < 0) o.timeWall = (strcmp(di->platformName, "Apple") == 0) ? 1 : 0; + if (o.timeWall) printf("timing: host wall time (Apple's OpenCL event timestamps are not usable; the rate is still a device rate, see README.md)\n"); + else printf("timing: device event profiling (CL_PROFILING_COMMAND_START/END), like cudaEvent elapsed time\n"); + + memset(&dv, 0, sizeof(dv)); + dv.ctx = clCreateContext(NULL, 1, &di->device, NULL, NULL, &err); + CL_CHECK_ERR(err, "clCreateContext"); + dv.q = clCreateCommandQueue(dv.ctx, di->device, CL_QUEUE_PROFILING_ENABLE, &err); + CL_CHECK_ERR(err, "clCreateCommandQueue"); + + src = readFile(o.kernelPath, &srcLen); + if (!src) { printf("FAIL: cannot read kernel source %s (run from proto-opencl/ or pass --kernel)\n", o.kernelPath); return 2; } + printf("kernel source: %s (%llu bytes)\n", o.kernelPath, (unsigned long long)srcLen); + setupProgram(&dv, di, &o, src, srcLen); + free(src); + printf("build options: %s\n", dv.buildOptions); + printf("exchange: %s\n", dv.exchangeNote); + { + size_t wg = 0; + cl_ulong lmem = 0; + clGetKernelWorkGroupInfo(dv.kHash, di->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(wg), &wg, NULL); + clGetKernelWorkGroupInfo(dv.kHash, di->device, CL_KERNEL_LOCAL_MEM_SIZE, sizeof(lmem), &lmem, NULL); + printf("kernel: igneum_hash max work-group %llu, local memory %llu bytes, work-group %d x 32\n", + (unsigned long long)wg, (unsigned long long)lmem, o.groupWarps); + } + printf("program: %d instructions x %d iterations, loads/hash %d, op mix %s\n", + IGNEUM_INSTR_COUNT, IGNEUM_ITERATIONS, IGNEUM_LOADS_PER_HASH, IGNEUM_OP_MIX); + printf("seed words: %08x %08x %08x %08x %08x %08x %08x %08x\n", + SEEDW[0], SEEDW[1], SEEDW[2], SEEDW[3], SEEDW[4], SEEDW[5], SEEDW[6], SEEDW[7]); + printf("day \"%s\" (d0 0x%08x, d1 0x%08x), pack dataset 2^%d words = %d MiB\n", + IGNEUM_DAY_STRING, IGNEUM_DAY0, IGNEUM_DAY1, IGNEUM_DATASET_LOG2, packMib()); +#if IGNEUM_DATASET_MODE == 1 + printf("dataset construction: memory-hard (256 MiB ChaCha cache, %d dependent cache reads per 64-byte item; proto-metal/MEMHARD.md)\n", IGNEUM_ITEM_ROUNDS); + cachePass = setupCache(&dv, di); +#else + printf("dataset construction: closed-form ds_elem (the original prototype dataset, not memory-hard)\n"); +#endif + + nonces = 1u << o.batchLog2; + if (nonces % (32u * (uint32_t)o.groupWarps) != 0u) { + printf("FAIL: 2^%d nonces is not a multiple of %d items per work-group\n", o.batchLog2, 32 * o.groupWarps); + return 2; + } + dOut = clCreateBuffer(dv.ctx, CL_MEM_READ_WRITE, (size_t)nonces * sizeof(uint64_t), NULL, &err); + CL_CHECK_ERR(err, "clCreateBuffer out"); + + if (o.sweep) { sizes[0] = 4; sizes[1] = 64; sizes[2] = 256; sizes[3] = 512; sizes[4] = 1024; nSizes = 5; } + else { sizes[0] = o.datasetMib; nSizes = 1; } + for (i = 0; i < nSizes; ++i) results[i] = runSize(&dv, di, &o, sizes[i], dOut, nonces); + CL_CHECK(clReleaseMemObject(dOut)); + + printf("\n=== summary (%s, %s, pack %s, batch 2^%d x %d, %d unit(s)/work-group, exchange %s, %s time) ===\n", + di->name, di->platformName, IGNEUM_SEED_STRING, o.batchLog2, o.batches, o.groupWarps, exchangeName(dv.exchange), o.timeWall ? "wall" : "device-event"); + printf("| dataset MiB | %s ms (second) | Mhash/s | GB/s useful | random loads/s (G) | loads/hash | dataset self-test | vectors |\n", + IGNEUM_DATASET_MODE == 1 ? "build" : "fill"); + printf("|---|---|---|---|---|---|---|---|\n"); + overall = cachePass; + for (i = 0; i < nSizes; ++i) { + const SizeResult* r = &results[i]; + overall = overall && r->dsPass && (!r->vecChecked || r->vecPass); + anyVec = anyVec || r->vecChecked; + printf("| %d | %.2f | %.3f | %.2f | %.2f | %d | %s | %s |\n", + r->mib, r->fillSecondMs, r->hashesPerSec / 1e6, r->gbps, + r->hashesPerSec * (double)IGNEUM_LOADS_PER_HASH / 1e9, IGNEUM_LOADS_PER_HASH, + r->dsPass ? "PASS" : "FAIL", + r->vecChecked ? (r->vecPass ? "PASS (3 warps, standalone and in batch)" : "FAIL") : "skipped (not pack size)"); + } +#if IGNEUM_DATASET_MODE == 1 + printf("cache: device fill %.2f ms (second), host fill %.1f ms one thread, cache check %s\n", gCacheFillSecondMs, gCacheHostMs, cachePass ? "PASS" : "FAIL"); + if (gCache) clReleaseMemObject(gCache); + free(hCache); +#endif + if (gProfilingFailures) printf("NOTE: %d event profiling queries failed on this runtime (timings printed as -1.00 ms are unavailable)\n", gProfilingFailures); + printf("exchange: %s\n", dv.exchangeNote); + if (!anyVec) printf("NOTE: no vectors were checked. Run at %d MiB (the default) to verify against the Mac.\n", packMib()); + printf("OVERALL: %s\n", overall ? "PASS" : "FAIL"); + + releaseProgram(&dv); + clReleaseCommandQueue(dv.q); + clReleaseContext(dv.ctx); + for (i = 0; i < nDev; ++i) free(devs[i].extensions); + free(devs); + return overall ? 0 : 1; +}