From c4e63c4419e404fcd840b56e9f811244bc46fdcb Mon Sep 17 00:00:00 2001 From: igneum-labs <337424239+igneum-labs@users.noreply.github.com> Date: Sat, 3 Oct 2026 16:02:05 +0000 Subject: [PATCH] Chain scene: checkpoints tagged from birth, locked tag, flash wash and final label Co-Authored-By: Claude Fable 5.1 --- proto-cuda/packs/igneum-genesis-mh/kernel.cu | 162 ++++++++++++++++++ proto-cuda/packs/igneum-genesis-mh/memhard.h | 101 +++++++++++ .../packs/igneum-genesis-mh/memhard.metal | 106 ++++++++++++ proto-cuda/packs/igneum-genesis-mh/program.h | 37 ++++ .../packs/igneum-genesis-mh/program.json | 114 ++++++++++++ .../packs/igneum-genesis-mh/program.metal | 108 ++++++++++++ proto-cuda/packs/igneum-genesis-mh/vectors.h | 53 ++++++ .../packs/igneum-genesis-mh/vectors.json | 36 ++++ site/index.html | 17 +- 9 files changed, 727 insertions(+), 7 deletions(-) create mode 100644 proto-cuda/packs/igneum-genesis-mh/kernel.cu create mode 100644 proto-cuda/packs/igneum-genesis-mh/memhard.h create mode 100644 proto-cuda/packs/igneum-genesis-mh/memhard.metal create mode 100644 proto-cuda/packs/igneum-genesis-mh/program.h create mode 100644 proto-cuda/packs/igneum-genesis-mh/program.json create mode 100644 proto-cuda/packs/igneum-genesis-mh/program.metal create mode 100644 proto-cuda/packs/igneum-genesis-mh/vectors.h create mode 100644 proto-cuda/packs/igneum-genesis-mh/vectors.json diff --git a/proto-cuda/packs/igneum-genesis-mh/kernel.cu b/proto-cuda/packs/igneum-genesis-mh/kernel.cu new file mode 100644 index 000000000..9fedfc99b --- /dev/null +++ b/proto-cuda/packs/igneum-genesis-mh/kernel.cu @@ -0,0 +1,162 @@ +// Generated by proto-metal/igneum-bench --export-pack for seed "igneum-genesis". Do not edit by hand. +// Bit-exact twin of the Metal kernel for the same seed (see proto-cuda/CHECKLIST.md and program.metal). +// Compiled ahead of time by nvcc together with proto-cuda/host.cu. No NVRTC. +#include +#include +#include "program.h" +#include "memhard.h" + +__device__ __forceinline__ uint32_t splitmix32(uint32_t 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, so both shift amounts are in 1..31. +__device__ __forceinline__ uint32_t rotl_imm(uint32_t x, uint32_t n) { return (x << n) | (x >> (32u - n)); } +// n is masked to 0..31; the second shift amount is masked too, so n == 0 gives x. +__device__ __forceinline__ uint32_t rotr_var(uint32_t x, uint32_t n) { n &= 31u; return (x >> n) | (x << ((32u - n) & 31u)); } +__device__ __forceinline__ uint32_t 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; +} + +// Memory-hard dataset (MEMHARD.md). One thread per cache segment; one thread per 64-byte dataset item. +// The core functions (mh_cache_segment, mh_item) are in memhard.h and are also compiled for the host. +__global__ void igneum_cache_fill(uint32_t* cache, uint32_t nSegments) { + uint32_t seg = blockIdx.x * blockDim.x + threadIdx.x; + if (seg < nSegments) mh_cache_segment(cache, seg); +} +__global__ void igneum_build(uint32_t* ds, const uint32_t* cache, uint32_t nItems) { + uint32_t t = blockIdx.x * blockDim.x + threadIdx.x; + if (t < nItems) { + uint32_t s[16]; + mh_item(cache, t, s); + uint32_t* d = ds + (size_t)t * 16u; + for (uint32_t i = 0u; i < 16u; ++i) d[i] = s[i]; + } +} +// One hash per thread. blockDim.x is a multiple of 32; lane = threadIdx.x & 31 and every +// __shfl_xor_sync stays inside the lane's own warp, exactly like simd_shuffle_xor inside a +// 32-wide Metal SIMD group. Control flow is uniform, so the full 0xffffffff member mask is valid. +__global__ void igneum_hash(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask) { + uint32_t gid = blockIdx.x * blockDim.x + threadIdx.x; + uint32_t nonce = baseNonce + gid; + uint32_t r0, r1, r2, r3, r4, r5, r6, r7; + { uint32_t x = nonce ^ 0x67a9a7beu; x += 0x9e3779b9u; x = splitmix32(x); r0 = x ^ 0x1a155b25u; } // SEEDW[0], 0x9e3779b9u * 1u, SEEDW[1] + { uint32_t x = nonce ^ 0x1a155b25u; x += 0x3c6ef372u; x = splitmix32(x); r1 = x ^ 0xfddfb732u; } // SEEDW[1], 0x9e3779b9u * 2u, SEEDW[2] + { uint32_t x = nonce ^ 0xfddfb732u; x += 0xdaa66d2bu; x = splitmix32(x); r2 = x ^ 0x4b5af2e8u; } // SEEDW[2], 0x9e3779b9u * 3u, SEEDW[3] + { uint32_t x = nonce ^ 0x4b5af2e8u; x += 0x78dde6e4u; x = splitmix32(x); r3 = x ^ 0xc55caf33u; } // SEEDW[3], 0x9e3779b9u * 4u, SEEDW[4] + { uint32_t x = nonce ^ 0xc55caf33u; x += 0x1715609du; x = splitmix32(x); r4 = x ^ 0xa27c13b7u; } // SEEDW[4], 0x9e3779b9u * 5u, SEEDW[5] + { uint32_t x = nonce ^ 0xa27c13b7u; x += 0xb54cda56u; x = splitmix32(x); r5 = x ^ 0x06628a48u; } // SEEDW[5], 0x9e3779b9u * 6u, SEEDW[6] + { uint32_t x = nonce ^ 0x06628a48u; x += 0x5384540fu; x = splitmix32(x); r6 = x ^ 0x03852469u; } // SEEDW[6], 0x9e3779b9u * 7u, SEEDW[7] + { uint32_t x = nonce ^ 0x03852469u; x += 0xf1bbcdc8u; x = splitmix32(x); r7 = x ^ 0x67a9a7beu; } // SEEDW[7], 0x9e3779b9u * 8u, SEEDW[0] + + for (uint32_t it = 0u; it < 8u; ++it) { + uint32_t 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 + r4 = r4 ^ __shfl_xor_sync(0xffffffffu, r6, 16); // 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 = __umulhi(r7, r4); // 22 mulhi + r5 = __umulhi(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 + r5 = r5 ^ __shfl_xor_sync(0xffffffffu, r1, 8); // 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 + r3 = r3 ^ __shfl_xor_sync(0xffffffffu, r0, 1); // 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 + r2 = r2 ^ __shfl_xor_sync(0xffffffffu, r3, 16); // 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 + r6 = r6 ^ __shfl_xor_sync(0xffffffffu, r2, 2); // 52 shfl + r2 = r2 - r0; // 53 sub + r0 = r0 ^ ds[r3 & mask]; // 54 load + r2 = r2 ^ __shfl_xor_sync(0xffffffffu, r1, 16); // 55 shfl + r0 = r0 + r6 + ((((sel >> 27u) & 1u) != 0u) ? 0xf4689674u : 0x25955401u); // 56 add + r2 = __umulhi(r2, r0); // 57 mulhi + r4 = __umulhi(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 = __umulhi(r7, r0); // 63 mulhi + } + uint32_t lo = r0 ^ rotl_imm(r1, 7u) ^ rotl_imm(r2, 14u) ^ rotl_imm(r3, 21u); + uint32_t hi = r4 ^ rotl_imm(r5, 9u) ^ rotl_imm(r6, 18u) ^ rotl_imm(r7, 27u); + out[gid] = ((uint64_t)hi << 32) | (uint64_t)lo; +} + +// Host-side launch wrappers. Declared in program.h, called from host.cu. +cudaError_t igneum_launch_cache_fill(uint32_t* cache, uint32_t nSegments) { + if (nSegments == 0u) return cudaErrorInvalidValue; + uint32_t block = 256u; + uint32_t grid = (nSegments + block - 1u) / block; + igneum_cache_fill<<>>(cache, nSegments); + return cudaGetLastError(); +} + +cudaError_t igneum_launch_build(uint32_t* ds, const uint32_t* cache, uint32_t nItems) { + if (nItems == 0u) return cudaErrorInvalidValue; + uint32_t block = 256u; + uint32_t grid = (nItems + block - 1u) / block; + igneum_build<<>>(ds, cache, nItems); + return cudaGetLastError(); +} +cudaError_t igneum_launch_hash(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask, + uint32_t nonces, uint32_t blockWarps) { + if (blockWarps == 0u || blockWarps > 32u) return cudaErrorInvalidValue; + uint32_t block = 32u * blockWarps; + if (nonces == 0u || (nonces % block) != 0u) return cudaErrorInvalidValue; + igneum_hash<<>>(ds, out, baseNonce, mask); + return cudaGetLastError(); +} + +cudaError_t igneum_hash_info(int* numRegs, int* blocksPerSM, uint32_t blockWarps) { + cudaFuncAttributes attr; + cudaError_t e = cudaFuncGetAttributes(&attr, igneum_hash); + if (e != cudaSuccess) return e; + *numRegs = attr.numRegs; + return cudaOccupancyMaxActiveBlocksPerMultiprocessor(blocksPerSM, igneum_hash, (int)(32u * blockWarps), 0); +} diff --git a/proto-cuda/packs/igneum-genesis-mh/memhard.h b/proto-cuda/packs/igneum-genesis-mh/memhard.h new file mode 100644 index 000000000..41aca0de6 --- /dev/null +++ b/proto-cuda/packs/igneum-genesis-mh/memhard.h @@ -0,0 +1,101 @@ +// 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. +#pragma once +#include +#if defined(__CUDACC__) +#define IGNEUM_HD __host__ __device__ __forceinline__ +#else +#define IGNEUM_HD static inline +#endif +// 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); } +IGNEUM_HD uint32_t mh_rotl(uint32_t x, uint32_t n) { return (x << n) | (x >> (32u - n)); } // n in 1..31 at every call site + +// y = ChaCha12 core(x) + x +IGNEUM_HD void mh_chacha_block(const uint32_t* x, uint32_t* y) { + for (uint32_t i = 0u; i < 16u; ++i) y[i] = x[i]; + for (uint32_t 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 (uint32_t 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. +IGNEUM_HD void mh_cache_segment(uint32_t* cache, uint32_t seg) { + uint32_t prev[16]; uint32_t x[16]; uint32_t y[16]; + for (uint32_t i = 0u; i < 16u; ++i) prev[i] = 0u; + for (uint32_t 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); + uint32_t* line = cache + ((seg * MH_SEGMENT_LINES + j) * 16u); + for (uint32_t 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. +IGNEUM_HD void mh_mixer(uint32_t* s, uint32_t 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. +IGNEUM_HD void mh_item(const uint32_t* cache, uint32_t t, uint32_t* 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 (uint32_t r = 0u; r < 8u; ++r) { + mh_mixer(s, 0x9E3779B9u * (r + 1u)); + const uint32_t* line = cache + ((s[0] & MH_CACHE_LINE_MASK) * 16u); + for (uint32_t 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. +IGNEUM_HD uint32_t mh_word(const uint32_t* cache, uint32_t w) { uint32_t s[16]; mh_item(cache, w >> 4u, s); return s[w & 15u]; } diff --git a/proto-cuda/packs/igneum-genesis-mh/memhard.metal b/proto-cuda/packs/igneum-genesis-mh/memhard.metal new file mode 100644 index 000000000..241866369 --- /dev/null +++ b/proto-cuda/packs/igneum-genesis-mh/memhard.metal @@ -0,0 +1,106 @@ +#include +using namespace metal; +// 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); } +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 +inline void mh_chacha_block(const thread uint* x, thread 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. +inline void mh_cache_segment(device 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); + device 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. +inline void mh_mixer(thread 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. +inline void mh_item(device const uint* cache, uint t, thread 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)); + device 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. +inline uint mh_word(device const uint* cache, uint w) { uint s[16]; mh_item(cache, w >> 4u, s); return s[w & 15u]; } + +// One thread per segment (2^16 threads). +kernel void igneum_cache_fill(device uint* cache [[buffer(0)]], uint gid [[thread_position_in_grid]]) { + mh_cache_segment(cache, gid); +} +// One thread per 64-byte item (dataset words / 16 threads). +kernel void igneum_build(device const uint* cache [[buffer(0)]], device uint* dataset [[buffer(1)]], + uint gid [[thread_position_in_grid]]) { + uint s[16]; + mh_item(cache, gid, s); + device uint* d = dataset + gid * 16u; + for (uint i = 0u; i < 16u; ++i) d[i] = s[i]; +} diff --git a/proto-cuda/packs/igneum-genesis-mh/program.h b/proto-cuda/packs/igneum-genesis-mh/program.h new file mode 100644 index 000000000..86c5e91c5 --- /dev/null +++ b/proto-cuda/packs/igneum-genesis-mh/program.h @@ -0,0 +1,37 @@ +// 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. +#pragma once +#include +#include + +#define IGNEUM_SEED_STRING "igneum-genesis" +#define IGNEUM_DAY_STRING "2026-10-03" +#define IGNEUM_DAY0 0x3067619fu +#define IGNEUM_DAY1 0x3c269176u +#define IGNEUM_DATASET_LOG2 28 +#define IGNEUM_MASK 0x0fffffffu +#define IGNEUM_LANES 32 +#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 1 + +#define IGNEUM_SEEDW_INIT { 0x67a9a7beu, 0x1a155b25u, 0xfddfb732u, 0x4b5af2e8u, 0xc55caf33u, 0xa27c13b7u, 0x06628a48u, 0x03852469u } +#define IGNEUM_KEY_INIT { 0x3067619fu, 0x3c269176u, 0x84a03b03u, 0xf8c63294u, 0xff977c5bu, 0xe60def3eu, 0x63630141u, 0xb8fbcb58u } +#define IGNEUM_CACHE_LOG2_WORDS 26 +#define IGNEUM_CACHE_SEGMENT_LOG2_LINES 6 +#define IGNEUM_CACHE_SEGMENTS 65536u +#define IGNEUM_ITEM_ROUNDS 8 +#define IGNEUM_MIX_ROT_INIT { 20u, 20u, 19u, 4u, 26u, 3u, 3u, 27u } +#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 } + +// 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); diff --git a/proto-cuda/packs/igneum-genesis-mh/program.json b/proto-cuda/packs/igneum-genesis-mh/program.json new file mode 100644 index 000000000..e28749a75 --- /dev/null +++ b/proto-cuda/packs/igneum-genesis-mh/program.json @@ -0,0 +1,114 @@ +{ + "format": "igneum-program-pack-2", + "dataset_mode": "memory-hard", + "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", + "lanes": 32, + "registers": 8, + "iterations": 8, + "instruction_count": 64, + "loads_per_hash": 104, + "op_mix": {"load": 13, "xor": 13, "sub": 7, "shfl": 6, "add": 5, "mulhi": 5, "mad": 4, "rotr": 4, "mul": 3, "rotl": 3, "or": 1}, + "register_init": "for i in 0..7: x = nonce ^ seed_words[i]; x += 0x9e3779b9 * (i+1) (mod 2^32); x = splitmix32(x); r[i] = x ^ seed_words[(i+1) & 7]", + "splitmix32": "x ^= x>>16; x *= 0x7feb352d; x ^= x>>15; x *= 0x846ca68b; x ^= x>>16", + "iteration": "sel = r0 sampled once at the top of each iteration, then all instructions in order", + "output": "lo = r0 ^ rotl(r1,7) ^ rotl(r2,14) ^ rotl(r3,21); hi = r4 ^ rotl(r5,9) ^ rotl(r6,18) ^ rotl(r7,27); out = (hi << 32) | lo", + "op_semantics": { + "add": "dst = dst + src + (bit `bit` of sel ? imm2 : imm)", + "sub": "dst = dst - src", + "mul": "dst = dst * src (low 32)", + "mulhi": "dst = high 32 bits of dst * src", + "xor": "dst = dst ^ src", + "or": "dst = dst | src", + "rotl": "dst = rotl(dst, rot), rot in 1..31", + "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]", + "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, + "bytes": 1073741824, + "mask": "0x0fffffff", + "day": "2026-10-03", + "day_words_from": "day/2026-10-03", + "d0": "0x3067619f", + "d1": "0x3c269176", + "mode": "memory-hard", + "spec": "proto-metal/MEMHARD.md", + "key": ["0x3067619f", "0x3c269176", "0x84a03b03", "0xf8c63294", "0xff977c5b", "0xe60def3e", "0x63630141", "0xb8fbcb58"], + "key_derivation": "the 8 words of seedWords(\"day/\" + day); d0, d1 are key[0], key[1]", + "cache": {"log2_words": 26, "bytes": 268435456, "line_words": 16, "segment_lines": 64, "segments": 65536, "block": "ChaCha12 core + feed-forward, rotations 16 12 8 7", "sigma": ["0x61707865", "0x3320646e", "0x79622d32", "0x6b206574"], "tag": ["0x49676e65", "0x756d4d48"], "chain": "in_j = prev_line ^ (sigma[0..3] || key[0..7] || seg || j || tag[0..1]); line_j = block(in_j); prev_0 = 0"}, + "mixer": {"draw": "SplitMix64 seeded with key[0] | key[1] << 32: rot[0..7] = 1 + next() % 31, mul[0..15] = low32(next()) | 1, rc[0..15] = low32(next())", "rot": [20, 20, 19, 4, 26, 3, 3, 27], "mul": ["0x42146205", "0x52cbe0fb", "0x7ecf4a03", "0x6728907f", "0xd81d9751", "0x132952c3", "0xf60de277", "0x05358035", "0xbaf6499d", "0xe4db9667", "0x3e98f45d", "0xd0004edd", "0x2691630d", "0x9beb3bcf", "0xab310379", "0x99cfb423"], "rc": ["0xbab68293", "0xcc162340", "0x6ce151cc", "0xe62b8997", "0xc9c80297", "0xf74a1654", "0x3d704af5", "0x3cf522b7", "0x2b9cac04", "0xa880ac10", "0x13e5dd1d", "0x6fc3e233", "0x2d83eeac", "0x9006e8bf", "0x2c4b5362", "0x31b49ee2"], "round": "for i in 0..15: s[i] = (s[i] ^ (rc[i] + (r+1) * 0x9E3779B9)) * mul[i]; then quarter rounds on columns (0,4,8,12) (1,5,9,13) (2,6,10,14) (3,7,11,15) with rot[0..3] and diagonals (0,5,10,15) (1,6,11,12) (2,7,8,13) (3,4,9,14) with rot[4..7]", "quarter_round": "a += b; d ^= a; d = rotl(d, r1); c += d; b ^= c; b = rotl(b, r2); a += b; d ^= a; d = rotl(d, r3); c += d; b ^= c; b = rotl(b, r4)"}, + "item": "s[0..7] = key; s[8+i] = t * mul[i] + rc[i] for i in 0..7; for r in 0..7: s = M_r(s); line = s[0] & "0x003fffff"; s[i] ^= cache[line * 16 + i]; then s = M_8(s); item(t) = s", + "word": "dataset[w] = item(w >> 4)[w & 15]" + }, + "instructions": [ + {"i": 0, "op": "rotl", "dst": 4, "src": 2, "src2": 7, "imm": "0x20699878", "imm2": "0x6f1a6170", "rot": 25, "bit": 7, "mask": 16}, + {"i": 1, "op": "sub", "dst": 0, "src": 5, "src2": 4, "imm": "0xf81a0b9d", "imm2": "0xf0505e88", "rot": 1, "bit": 4, "mask": 1}, + {"i": 2, "op": "load", "dst": 4, "src": 3, "src2": 6, "imm": "0xb1978a0b", "imm2": "0x2ca4e162", "rot": 10, "bit": 21, "mask": 1}, + {"i": 3, "op": "rotl", "dst": 1, "src": 6, "src2": 7, "imm": "0xc3bd2355", "imm2": "0xa8c5f27e", "rot": 1, "bit": 22, "mask": 1}, + {"i": 4, "op": "add", "dst": 2, "src": 3, "src2": 0, "imm": "0x61f0b51c", "imm2": "0x2735a174", "rot": 4, "bit": 26, "mask": 2}, + {"i": 5, "op": "load", "dst": 5, "src": 3, "src2": 3, "imm": "0x4d183796", "imm2": "0x679648a8", "rot": 4, "bit": 30, "mask": 4}, + {"i": 6, "op": "sub", "dst": 5, "src": 7, "src2": 7, "imm": "0x265677dc", "imm2": "0x9043323e", "rot": 30, "bit": 7, "mask": 4}, + {"i": 7, "op": "add", "dst": 3, "src": 4, "src2": 6, "imm": "0x52f2dbf4", "imm2": "0x5a069596", "rot": 5, "bit": 26, "mask": 2}, + {"i": 8, "op": "xor", "dst": 0, "src": 4, "src2": 1, "imm": "0x3303ec4b", "imm2": "0xfeca75be", "rot": 21, "bit": 7, "mask": 1}, + {"i": 9, "op": "xor", "dst": 4, "src": 0, "src2": 0, "imm": "0x5dc5959c", "imm2": "0x023f44a8", "rot": 11, "bit": 7, "mask": 4}, + {"i": 10, "op": "load", "dst": 2, "src": 0, "src2": 3, "imm": "0xde716173", "imm2": "0xc21e924d", "rot": 30, "bit": 15, "mask": 1}, + {"i": 11, "op": "shfl", "dst": 4, "src": 6, "src2": 1, "imm": "0x3cda6d48", "imm2": "0x0970145b", "rot": 12, "bit": 22, "mask": 16}, + {"i": 12, "op": "sub", "dst": 1, "src": 5, "src2": 4, "imm": "0x48866c15", "imm2": "0x4eaee50c", "rot": 30, "bit": 20, "mask": 16}, + {"i": 13, "op": "xor", "dst": 2, "src": 1, "src2": 5, "imm": "0x196d165c", "imm2": "0x0f730511", "rot": 11, "bit": 3, "mask": 4}, + {"i": 14, "op": "load", "dst": 4, "src": 5, "src2": 1, "imm": "0xc5c3b55d", "imm2": "0xec061424", "rot": 26, "bit": 27, "mask": 8}, + {"i": 15, "op": "load", "dst": 2, "src": 4, "src2": 1, "imm": "0x17c9c95b", "imm2": "0x306542fe", "rot": 27, "bit": 17, "mask": 16}, + {"i": 16, "op": "load", "dst": 3, "src": 0, "src2": 2, "imm": "0x590f9e11", "imm2": "0xa4c13036", "rot": 3, "bit": 28, "mask": 16}, + {"i": 17, "op": "xor", "dst": 4, "src": 6, "src2": 3, "imm": "0x9c4eeee9", "imm2": "0xf069b834", "rot": 11, "bit": 23, "mask": 8}, + {"i": 18, "op": "mad", "dst": 2, "src": 4, "src2": 6, "imm": "0xa242a28b", "imm2": "0xb8974bdf", "rot": 13, "bit": 30, "mask": 1}, + {"i": 19, "op": "rotr", "dst": 6, "src": 1, "src2": 0, "imm": "0xcd69ed50", "imm2": "0xadf52615", "rot": 30, "bit": 2, "mask": 2}, + {"i": 20, "op": "xor", "dst": 3, "src": 4, "src2": 0, "imm": "0xe3059a24", "imm2": "0x16b8dd86", "rot": 25, "bit": 31, "mask": 16}, + {"i": 21, "op": "mad", "dst": 1, "src": 3, "src2": 5, "imm": "0xc3ae8ae1", "imm2": "0x8e126e8a", "rot": 2, "bit": 28, "mask": 1}, + {"i": 22, "op": "mulhi", "dst": 7, "src": 4, "src2": 2, "imm": "0x6909af7a", "imm2": "0xa388b4b9", "rot": 14, "bit": 10, "mask": 2}, + {"i": 23, "op": "mulhi", "dst": 5, "src": 2, "src2": 2, "imm": "0xdf099cfb", "imm2": "0xd0133a01", "rot": 31, "bit": 3, "mask": 4}, + {"i": 24, "op": "load", "dst": 0, "src": 6, "src2": 3, "imm": "0x0a3056de", "imm2": "0x7f0c25c3", "rot": 27, "bit": 13, "mask": 8}, + {"i": 25, "op": "mul", "dst": 5, "src": 6, "src2": 7, "imm": "0x089f5404", "imm2": "0xbd066e1d", "rot": 7, "bit": 10, "mask": 4}, + {"i": 26, "op": "load", "dst": 7, "src": 1, "src2": 2, "imm": "0x3e26afea", "imm2": "0xba573970", "rot": 11, "bit": 15, "mask": 8}, + {"i": 27, "op": "rotr", "dst": 3, "src": 1, "src2": 5, "imm": "0xe2466d63", "imm2": "0x30db112b", "rot": 31, "bit": 15, "mask": 2}, + {"i": 28, "op": "shfl", "dst": 5, "src": 1, "src2": 3, "imm": "0xdc20297d", "imm2": "0xcb813557", "rot": 3, "bit": 21, "mask": 8}, + {"i": 29, "op": "xor", "dst": 7, "src": 5, "src2": 2, "imm": "0x6ede5c13", "imm2": "0x1f14267f", "rot": 22, "bit": 31, "mask": 1}, + {"i": 30, "op": "rotl", "dst": 7, "src": 5, "src2": 6, "imm": "0x24cdb54b", "imm2": "0xa44e008d", "rot": 23, "bit": 4, "mask": 16}, + {"i": 31, "op": "sub", "dst": 2, "src": 0, "src2": 2, "imm": "0xf94d9b65", "imm2": "0x3457264c", "rot": 11, "bit": 5, "mask": 2}, + {"i": 32, "op": "xor", "dst": 7, "src": 2, "src2": 3, "imm": "0x8d2046e5", "imm2": "0xf68893bb", "rot": 21, "bit": 2, "mask": 1}, + {"i": 33, "op": "xor", "dst": 2, "src": 6, "src2": 5, "imm": "0xe0ebc4ce", "imm2": "0x02773069", "rot": 21, "bit": 16, "mask": 1}, + {"i": 34, "op": "load", "dst": 6, "src": 1, "src2": 2, "imm": "0x3b2d2124", "imm2": "0x187a9128", "rot": 1, "bit": 9, "mask": 16}, + {"i": 35, "op": "xor", "dst": 1, "src": 4, "src2": 6, "imm": "0xe3f24158", "imm2": "0x5c64a589", "rot": 13, "bit": 0, "mask": 8}, + {"i": 36, "op": "shfl", "dst": 3, "src": 0, "src2": 3, "imm": "0x63578bc1", "imm2": "0xbf64a89f", "rot": 16, "bit": 26, "mask": 1}, + {"i": 37, "op": "mul", "dst": 2, "src": 6, "src2": 4, "imm": "0xfa624b69", "imm2": "0x0389cf85", "rot": 28, "bit": 4, "mask": 1}, + {"i": 38, "op": "add", "dst": 5, "src": 3, "src2": 5, "imm": "0x71f30417", "imm2": "0xa29f4338", "rot": 2, "bit": 12, "mask": 2}, + {"i": 39, "op": "xor", "dst": 7, "src": 6, "src2": 3, "imm": "0xfa3b845c", "imm2": "0x9717695d", "rot": 18, "bit": 24, "mask": 1}, + {"i": 40, "op": "xor", "dst": 7, "src": 3, "src2": 4, "imm": "0x8408dc67", "imm2": "0x8d389f9b", "rot": 21, "bit": 16, "mask": 16}, + {"i": 41, "op": "rotr", "dst": 3, "src": 4, "src2": 1, "imm": "0x4bbcad92", "imm2": "0xbc6375cd", "rot": 18, "bit": 4, "mask": 2}, + {"i": 42, "op": "xor", "dst": 5, "src": 3, "src2": 3, "imm": "0xbd31a2ea", "imm2": "0x742d28ea", "rot": 4, "bit": 12, "mask": 4}, + {"i": 43, "op": "rotr", "dst": 3, "src": 6, "src2": 5, "imm": "0x35c52d04", "imm2": "0x0f8e8621", "rot": 3, "bit": 22, "mask": 8}, + {"i": 44, "op": "mad", "dst": 1, "src": 3, "src2": 5, "imm": "0x3958f280", "imm2": "0x8713c7e1", "rot": 5, "bit": 23, "mask": 16}, + {"i": 45, "op": "add", "dst": 7, "src": 3, "src2": 7, "imm": "0xc1ae8d3b", "imm2": "0xa907b90b", "rot": 13, "bit": 29, "mask": 1}, + {"i": 46, "op": "or", "dst": 7, "src": 2, "src2": 2, "imm": "0xc4f18bec", "imm2": "0x8a3e2464", "rot": 30, "bit": 16, "mask": 2}, + {"i": 47, "op": "shfl", "dst": 2, "src": 3, "src2": 6, "imm": "0xf73ba9e3", "imm2": "0x028aa63c", "rot": 3, "bit": 20, "mask": 16}, + {"i": 48, "op": "load", "dst": 1, "src": 7, "src2": 5, "imm": "0x53be87d4", "imm2": "0x690a5729", "rot": 1, "bit": 17, "mask": 8}, + {"i": 49, "op": "sub", "dst": 5, "src": 1, "src2": 7, "imm": "0xecdc31a6", "imm2": "0xed6bd9f2", "rot": 1, "bit": 13, "mask": 8}, + {"i": 50, "op": "load", "dst": 3, "src": 1, "src2": 3, "imm": "0x55f91dbc", "imm2": "0xf026553d", "rot": 30, "bit": 27, "mask": 1}, + {"i": 51, "op": "sub", "dst": 2, "src": 3, "src2": 5, "imm": "0x6689eef2", "imm2": "0x24f3ae21", "rot": 6, "bit": 22, "mask": 8}, + {"i": 52, "op": "shfl", "dst": 6, "src": 2, "src2": 4, "imm": "0x8fd37aad", "imm2": "0x89ed5d67", "rot": 8, "bit": 16, "mask": 2}, + {"i": 53, "op": "sub", "dst": 2, "src": 0, "src2": 0, "imm": "0x963bb7e6", "imm2": "0x4738b84f", "rot": 18, "bit": 16, "mask": 4}, + {"i": 54, "op": "load", "dst": 0, "src": 3, "src2": 0, "imm": "0x838b5065", "imm2": "0x36360066", "rot": 3, "bit": 31, "mask": 4}, + {"i": 55, "op": "shfl", "dst": 2, "src": 1, "src2": 5, "imm": "0xb17aad78", "imm2": "0x8458f7ac", "rot": 5, "bit": 7, "mask": 16}, + {"i": 56, "op": "add", "dst": 0, "src": 6, "src2": 0, "imm": "0x25955401", "imm2": "0xf4689674", "rot": 14, "bit": 27, "mask": 4}, + {"i": 57, "op": "mulhi", "dst": 2, "src": 0, "src2": 0, "imm": "0xa4e8c86a", "imm2": "0x14bde8e1", "rot": 8, "bit": 9, "mask": 16}, + {"i": 58, "op": "mulhi", "dst": 4, "src": 2, "src2": 7, "imm": "0xa36b1f4f", "imm2": "0x372e072f", "rot": 31, "bit": 31, "mask": 1}, + {"i": 59, "op": "mad", "dst": 2, "src": 6, "src2": 7, "imm": "0xfcd3b17f", "imm2": "0x6a8583df", "rot": 24, "bit": 9, "mask": 2}, + {"i": 60, "op": "xor", "dst": 3, "src": 1, "src2": 0, "imm": "0x3bbea2ae", "imm2": "0x914e2f0a", "rot": 2, "bit": 15, "mask": 16}, + {"i": 61, "op": "mul", "dst": 4, "src": 2, "src2": 6, "imm": "0x50aaa099", "imm2": "0xd1b988a1", "rot": 27, "bit": 31, "mask": 4}, + {"i": 62, "op": "load", "dst": 0, "src": 2, "src2": 3, "imm": "0x4ccdbdea", "imm2": "0xf3a60b41", "rot": 1, "bit": 20, "mask": 2}, + {"i": 63, "op": "mulhi", "dst": 7, "src": 0, "src2": 7, "imm": "0x3e400372", "imm2": "0x92c6201f", "rot": 1, "bit": 13, "mask": 8} + ] +} diff --git a/proto-cuda/packs/igneum-genesis-mh/program.metal b/proto-cuda/packs/igneum-genesis-mh/program.metal new file mode 100644 index 000000000..65910c0e7 --- /dev/null +++ b/proto-cuda/packs/igneum-genesis-mh/program.metal @@ -0,0 +1,108 @@ +#include +using namespace metal; + +#define MASK 0x0fffffffu +constant uint SEEDW[8] = { 0x67a9a7beu, 0x1a155b25u, 0xfddfb732u, 0x4b5af2e8u, 0xc55caf33u, 0xa27c13b7u, 0x06628a48u, 0x03852469u }; + +inline uint splitmix32(uint x) { + x ^= x >> 16; x *= 0x7feb352du; + x ^= x >> 15; x *= 0x846ca68bu; + x ^= x >> 16; + return x; +} +inline uint rotl_imm(uint x, uint n) { return (x << n) | (x >> (32u - n)); } // n in 1..31 +inline uint rotr_var(uint x, uint n) { n &= 31u; return (x >> n) | (x << ((32u - n) & 31u)); } +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; +} +kernel void igneum_hash(device const uint* dataset [[buffer(0)]], + device ulong* out [[buffer(1)]], + constant uint& baseNonce [[buffer(2)]], + uint gid [[thread_position_in_grid]]) { + uint nonce = baseNonce + gid; + uint r0, r1, r2, r3, r4, r5, r6, r7; + { uint x = nonce ^ SEEDW[0]; x += 0x9e3779b9u * 1u; x = splitmix32(x); r0 = x ^ SEEDW[1]; } + { uint x = nonce ^ SEEDW[1]; x += 0x9e3779b9u * 2u; x = splitmix32(x); r1 = x ^ SEEDW[2]; } + { uint x = nonce ^ SEEDW[2]; x += 0x9e3779b9u * 3u; x = splitmix32(x); r2 = x ^ SEEDW[3]; } + { uint x = nonce ^ SEEDW[3]; x += 0x9e3779b9u * 4u; x = splitmix32(x); r3 = x ^ SEEDW[4]; } + { uint x = nonce ^ SEEDW[4]; x += 0x9e3779b9u * 5u; x = splitmix32(x); r4 = x ^ SEEDW[5]; } + { uint x = nonce ^ SEEDW[5]; x += 0x9e3779b9u * 6u; x = splitmix32(x); r5 = x ^ SEEDW[6]; } + { uint x = nonce ^ SEEDW[6]; x += 0x9e3779b9u * 7u; x = splitmix32(x); r6 = x ^ SEEDW[7]; } + { uint x = nonce ^ SEEDW[7]; x += 0x9e3779b9u * 8u; x = splitmix32(x); r7 = x ^ SEEDW[0]; } + + for (uint it = 0u; it < 8u; ++it) { + uint sel = r0; + r4 = rotl_imm(r4, 25u); // 0 + r0 = r0 - r5; // 1 + r4 = r4 ^ dataset[r3 & MASK]; // 2 + r1 = rotl_imm(r1, 1u); // 3 + r2 = r2 + r3 + select(0x61f0b51cu, 0x2735a174u, ((sel >> 26u) & 1u) != 0u); // 4 + r5 = r5 ^ dataset[r3 & MASK]; // 5 + r5 = r5 - r7; // 6 + r3 = r3 + r4 + select(0x52f2dbf4u, 0x5a069596u, ((sel >> 26u) & 1u) != 0u); // 7 + r0 = r0 ^ r4; // 8 + r4 = r4 ^ r0; // 9 + r2 = r2 ^ dataset[r0 & MASK]; // 10 + r4 = r4 ^ simd_shuffle_xor(r6, (ushort)16); // 11 + r1 = r1 - r5; // 12 + r2 = r2 ^ r1; // 13 + r4 = r4 ^ dataset[r5 & MASK]; // 14 + r2 = r2 ^ dataset[r4 & MASK]; // 15 + r3 = r3 ^ dataset[r0 & MASK]; // 16 + r4 = r4 ^ r6; // 17 + r2 = r4 * r6 + r2; // 18 + r6 = rotr_var(r6, r1); // 19 + r3 = r3 ^ r4; // 20 + r1 = r3 * r5 + r1; // 21 + r7 = mulhi(r7, r4); // 22 + r5 = mulhi(r5, r2); // 23 + r0 = r0 ^ dataset[r6 & MASK]; // 24 + r5 = r5 * r6; // 25 + r7 = r7 ^ dataset[r1 & MASK]; // 26 + r3 = rotr_var(r3, r1); // 27 + r5 = r5 ^ simd_shuffle_xor(r1, (ushort)8); // 28 + r7 = r7 ^ r5; // 29 + r7 = rotl_imm(r7, 23u); // 30 + r2 = r2 - r0; // 31 + r7 = r7 ^ r2; // 32 + r2 = r2 ^ r6; // 33 + r6 = r6 ^ dataset[r1 & MASK]; // 34 + r1 = r1 ^ r4; // 35 + r3 = r3 ^ simd_shuffle_xor(r0, (ushort)1); // 36 + r2 = r2 * r6; // 37 + r5 = r5 + r3 + select(0x71f30417u, 0xa29f4338u, ((sel >> 12u) & 1u) != 0u); // 38 + r7 = r7 ^ r6; // 39 + r7 = r7 ^ r3; // 40 + r3 = rotr_var(r3, r4); // 41 + r5 = r5 ^ r3; // 42 + r3 = rotr_var(r3, r6); // 43 + r1 = r3 * r5 + r1; // 44 + r7 = r7 + r3 + select(0xc1ae8d3bu, 0xa907b90bu, ((sel >> 29u) & 1u) != 0u); // 45 + r7 = r7 | r2; // 46 + r2 = r2 ^ simd_shuffle_xor(r3, (ushort)16); // 47 + r1 = r1 ^ dataset[r7 & MASK]; // 48 + r5 = r5 - r1; // 49 + r3 = r3 ^ dataset[r1 & MASK]; // 50 + r2 = r2 - r3; // 51 + r6 = r6 ^ simd_shuffle_xor(r2, (ushort)2); // 52 + r2 = r2 - r0; // 53 + r0 = r0 ^ dataset[r3 & MASK]; // 54 + r2 = r2 ^ simd_shuffle_xor(r1, (ushort)16); // 55 + r0 = r0 + r6 + select(0x25955401u, 0xf4689674u, ((sel >> 27u) & 1u) != 0u); // 56 + r2 = mulhi(r2, r0); // 57 + r4 = mulhi(r4, r2); // 58 + r2 = r6 * r7 + r2; // 59 + r3 = r3 ^ r1; // 60 + r4 = r4 * r2; // 61 + r0 = r0 ^ dataset[r2 & MASK]; // 62 + r7 = mulhi(r7, r0); // 63 + } + 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; +} diff --git a/proto-cuda/packs/igneum-genesis-mh/vectors.h b/proto-cuda/packs/igneum-genesis-mh/vectors.h new file mode 100644 index 000000000..2c97eb016 --- /dev/null +++ b/proto-cuda/packs/igneum-genesis-mh/vectors.h @@ -0,0 +1,53 @@ +// 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 +#include + +#define IGNEUM_VEC_WARPS 3 +static const uint32_t IGNEUM_VEC_BASE[IGNEUM_VEC_WARPS] = { 0u, 4096u, 1000000u }; +static const uint64_t IGNEUM_VEC_OUT[IGNEUM_VEC_WARPS][32] = { + { // base nonce 0 + 0x1fb0b3bbc1ac8279ull, 0x61533759ca995ac8ull, 0x97cd15ec31242001ull, 0x3c19e71e2dbff828ull, 0xc7aba60b6c2ab016ull, 0x8f13725d2a59bf84ull, 0x7b826b82fdec5a2full, 0x3131d7817418c04eull, + 0xd7c7e57c6aaac948ull, 0xb3b639431ebdd3d6ull, 0x400af448658a1d56ull, 0x17f5760b8cadb8fdull, 0x01ded5c4894c411cull, 0xd6b3fdbc57bdb128ull, 0x888efb103cecb983ull, 0xac73538353a340c4ull, + 0x3f9914ad9445f052ull, 0x8a860b28afd3dd42ull, 0xc2a122ef9c1fa330ull, 0x96c4e58663f82dceull, 0xbaae4b9d6a3c4320ull, 0xcd6c2dd653e07743ull, 0x7d3d00f0fb46325bull, 0x7636224050132b09ull, + 0xdc5d76c741cc60f4ull, 0xaef3031b0d4b701cull, 0x735987104a975f5full, 0xa58fce4010bf99ddull, 0xb16c863d2ffdc0ddull, 0x0337b56e39529c90ull, 0x802ac0bc1d0c0696ull, 0xfa052263a854f3deull + }, + { // base nonce 4096 + 0xf219cf7ecf6ec450ull, 0xb575ab4b388c01f7ull, 0x09ff62fd27055f6dull, 0x0c8abc61759d3564ull, 0xa5909acbcc981201ull, 0x1272880e65dc626dull, 0xf49baed726dabd3aull, 0xd9bf25b9d7cd2c59ull, + 0xc1a453f0e2a0e26eull, 0x9c8bd1806912812full, 0xc7bc5603edeead7eull, 0x9a0218abf9fa1fd0ull, 0x01affc92769181b7ull, 0xe54fc07ff716c78aull, 0x70606c49b68b2fb3ull, 0x51a7e9c37930bcc3ull, + 0x47540e9d54ce6117ull, 0xd1dfe07ffafb6952ull, 0x701a455888ed0a2aull, 0x4baec8d5ffb99fadull, 0xc6dffc6807b38113ull, 0xdc9c3cabc5965b78ull, 0xa3e4e56eee3fe8d6ull, 0xf6decd4c32d0b018ull, + 0x85cd622043eb59b1ull, 0x4642d5079564dcdfull, 0x84e19ac8eefef8acull, 0x2f4414ad779564deull, 0xa288c93ca1c51b26ull, 0xa977d793f289fe1cull, 0xa1bc6f71f19384ffull, 0xd5b198435d9c40daull + }, + { // base nonce 1000000 + 0x8a6f7a32b06edb52ull, 0x362d2b789a0c0ebfull, 0x944ec8bd1d08c026ull, 0x6c40b67428eb7af3ull, 0x306e546583cb011eull, 0x27a55ab280344f21ull, 0xa6247ff1d8e72bceull, 0xf11780aa78365f00ull, + 0x2919860ba64f212aull, 0x5cfa0a33f4dd691dull, 0xce626a9cc0d2427dull, 0xf42f1047cdfbcb8bull, 0xd247b8ab5c35aeaaull, 0x47989e53b0251facull, 0xb55e4e5fa8df8217ull, 0x66b17338a66756caull, + 0x34b56d64da199729ull, 0x167a7c11aad8f00dull, 0x4292a627e69d80cbull, 0xce30b0820c123a4bull, 0xcd3b7367ba86adadull, 0xb0c1b22a10274c28ull, 0x7326d56bc7140e08ull, 0xea6490cba6bb2b4full, + 0xab57b3e93a1dae5cull, 0x2f3c7c5a624a8f0aull, 0x1f9a062062647b0bull, 0x83cf530898ba80f8ull, 0xa9f28306c97622baull, 0xe4212399129bddeeull, 0x8a2e1f8094e0a1f4ull, 0x7040299873672a35ull + } +}; + +// Dataset self-test: dataset[0..15] and dataset[IGNEUM_MASK] (268435455). +static const uint32_t IGNEUM_DS_HEAD[16] = { + 0xffc3cd94u, 0x5920ccd8u, 0x392f44bbu, 0x5e57f67au, 0x2f2bc2a9u, 0x620b0e36u, 0xbdc09014u, 0x436654bfu, + 0x311e0b48u, 0x1abd93adu, 0x59cc7ce8u, 0xee5247b2u, 0x86171fe8u, 0x6d874751u, 0xc9f7728fu, 0x7c2a435du +}; +static const uint32_t IGNEUM_DS_LAST_INDEX = 268435455u; +static const uint32_t IGNEUM_DS_LAST = 0xa33ada72u; +// 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] = { + 0xe8b73d94u, 0x337028b5u, 0xafe148c9u, 0xab99f7aeu, 0x434ea619u, 0xd85cb880u, 0x54764c7fu, 0x82c7e420u, 0xedf4cb9eu, 0x9884c959u, 0x223ee793u, 0x3a9ccf69u, 0x81da4fd2u, 0xd6ce8cb9u, 0xe3922dcau, 0x3e7e6bdeu, 0x382a3acau, 0x567e7f7fu, 0x25a0f084u, 0xbfeef128u, 0xe338abfbu, 0x7c3b5280u, 0x909bc5f1u, 0xd8b74b9cu, 0x8e31a22eu, 0x26b5f1d8u, 0x79122c00u, 0xcafc3340u, 0xd5e02ea3u, 0x1aee1afdu, 0xdb090d9au, 0xb049f435u, 0x4954d8bau, 0x03797ba0u, 0x196eefbdu, 0xd153412au, 0xbe5d2c4bu, 0xdaa14f0eu, 0x8e61ed07u, 0x9e9a64c6u, 0x2e29ff36u, 0x392a8589u, 0xb56a5912u, 0xfa6e8b57u, 0xd1a737cbu, 0xb0fa841au, 0xbe1c341fu, 0xe25be0f1u, 0xe937f543u, 0xebab2248u, 0x8e1b607au, 0x202a2fedu, 0x95e2819cu, 0x9c9652d4u, 0x32fedef0u, 0xdecfff82u, 0xcb5d43e5u, 0xb735806au, 0x8905939cu, 0xfbf8472du, 0xada74e5du, 0x7ebdeeeau, 0x0119f2b3u, 0xa9a376b8u +}; +// Cache self-test (memory-hard mode): cache[0..15], the last 16 words, and FNV-1a 64 over all 2^26 words. +static const uint32_t IGNEUM_CACHE_HEAD[16] = { + 0x355a86d2u, 0x7957db1cu, 0xd21772afu, 0x6fc1e09bu, 0xd55ce61du, 0x6e6a278bu, 0xd3f543ceu, 0x223d8e82u, + 0x143ab337u, 0x2e9f05bdu, 0x2eb389bfu, 0x0c6e449eu, 0x5cfa4222u, 0xba6560feu, 0x8e3e1aa4u, 0xdbcc1d53u +}; +static const uint32_t IGNEUM_CACHE_LAST[16] = { + 0x41190d91u, 0xbd277957u, 0x22ddbb49u, 0x6986f207u, 0xdf69a4d6u, 0x26401a3au, 0x818230fbu, 0xc417122du, + 0x3597b211u, 0xb553ce55u, 0xcf39cc0du, 0x3b7fc43au, 0x3fd43b00u, 0x67e1c80eu, 0xffa7ea7du, 0xca2960abu +}; +static const uint64_t IGNEUM_CACHE_FNV64 = 0x48c4f5bf24166b2eull; diff --git a/proto-cuda/packs/igneum-genesis-mh/vectors.json b/proto-cuda/packs/igneum-genesis-mh/vectors.json new file mode 100644 index 000000000..507a4678b --- /dev/null +++ b/proto-cuda/packs/igneum-genesis-mh/vectors.json @@ -0,0 +1,36 @@ +{ + "seed": "igneum-genesis", + "day": "2026-10-03", + "dataset_mode": "memory-hard", + "dataset_log2_words": 28, + "mask": "0x0fffffff", + "lanes": 32, + "source": "proto-metal CPU interpreter (cpuWarp, memory-hard dataset) on Apple M5 Max; Metal GPU cross-check PASS 3/3 warps", + "warps": [ + {"base_nonce": 0, "expected": [ + "0x1fb0b3bbc1ac8279", "0x61533759ca995ac8", "0x97cd15ec31242001", "0x3c19e71e2dbff828", "0xc7aba60b6c2ab016", "0x8f13725d2a59bf84", "0x7b826b82fdec5a2f", "0x3131d7817418c04e", + "0xd7c7e57c6aaac948", "0xb3b639431ebdd3d6", "0x400af448658a1d56", "0x17f5760b8cadb8fd", "0x01ded5c4894c411c", "0xd6b3fdbc57bdb128", "0x888efb103cecb983", "0xac73538353a340c4", + "0x3f9914ad9445f052", "0x8a860b28afd3dd42", "0xc2a122ef9c1fa330", "0x96c4e58663f82dce", "0xbaae4b9d6a3c4320", "0xcd6c2dd653e07743", "0x7d3d00f0fb46325b", "0x7636224050132b09", + "0xdc5d76c741cc60f4", "0xaef3031b0d4b701c", "0x735987104a975f5f", "0xa58fce4010bf99dd", "0xb16c863d2ffdc0dd", "0x0337b56e39529c90", "0x802ac0bc1d0c0696", "0xfa052263a854f3de" + ]}, + {"base_nonce": 4096, "expected": [ + "0xf219cf7ecf6ec450", "0xb575ab4b388c01f7", "0x09ff62fd27055f6d", "0x0c8abc61759d3564", "0xa5909acbcc981201", "0x1272880e65dc626d", "0xf49baed726dabd3a", "0xd9bf25b9d7cd2c59", + "0xc1a453f0e2a0e26e", "0x9c8bd1806912812f", "0xc7bc5603edeead7e", "0x9a0218abf9fa1fd0", "0x01affc92769181b7", "0xe54fc07ff716c78a", "0x70606c49b68b2fb3", "0x51a7e9c37930bcc3", + "0x47540e9d54ce6117", "0xd1dfe07ffafb6952", "0x701a455888ed0a2a", "0x4baec8d5ffb99fad", "0xc6dffc6807b38113", "0xdc9c3cabc5965b78", "0xa3e4e56eee3fe8d6", "0xf6decd4c32d0b018", + "0x85cd622043eb59b1", "0x4642d5079564dcdf", "0x84e19ac8eefef8ac", "0x2f4414ad779564de", "0xa288c93ca1c51b26", "0xa977d793f289fe1c", "0xa1bc6f71f19384ff", "0xd5b198435d9c40da" + ]}, + {"base_nonce": 1000000, "expected": [ + "0x8a6f7a32b06edb52", "0x362d2b789a0c0ebf", "0x944ec8bd1d08c026", "0x6c40b67428eb7af3", "0x306e546583cb011e", "0x27a55ab280344f21", "0xa6247ff1d8e72bce", "0xf11780aa78365f00", + "0x2919860ba64f212a", "0x5cfa0a33f4dd691d", "0xce626a9cc0d2427d", "0xf42f1047cdfbcb8b", "0xd247b8ab5c35aeaa", "0x47989e53b0251fac", "0xb55e4e5fa8df8217", "0x66b17338a66756ca", + "0x34b56d64da199729", "0x167a7c11aad8f00d", "0x4292a627e69d80cb", "0xce30b0820c123a4b", "0xcd3b7367ba86adad", "0xb0c1b22a10274c28", "0x7326d56bc7140e08", "0xea6490cba6bb2b4f", + "0xab57b3e93a1dae5c", "0x2f3c7c5a624a8f0a", "0x1f9a062062647b0b", "0x83cf530898ba80f8", "0xa9f28306c97622ba", "0xe4212399129bddee", "0x8a2e1f8094e0a1f4", "0x7040299873672a35" + ]} + ], + "dataset_head": ["0xffc3cd94", "0x5920ccd8", "0x392f44bb", "0x5e57f67a", "0x2f2bc2a9", "0x620b0e36", "0xbdc09014", "0x436654bf", "0x311e0b48", "0x1abd93ad", "0x59cc7ce8", "0xee5247b2", "0x86171fe8", "0x6d874751", "0xc9f7728f", "0x7c2a435d"], + "dataset_last_index": 268435455, + "dataset_last": "0xa33ada72", + "dataset_samples": [{"index": 59471966, "value": "0xe8b73d94"}, {"index": 217795994, "value": "0x337028b5"}, {"index": 208353206, "value": "0xafe148c9"}, {"index": 42483309, "value": "0xab99f7ae"}, {"index": 172547758, "value": "0x434ea619"}, {"index": 148076330, "value": "0xd85cb880"}, {"index": 183853158, "value": "0x54764c7f"}, {"index": 214389424, "value": "0x82c7e420"}, {"index": 267488061, "value": "0xedf4cb9e"}, {"index": 169781097, "value": "0x9884c959"}, {"index": 184093494, "value": "0x223ee793"}, {"index": 153880993, "value": "0x3a9ccf69"}, {"index": 84977930, "value": "0x81da4fd2"}, {"index": 46426879, "value": "0xd6ce8cb9"}, {"index": 3093825, "value": "0xe3922dca"}, {"index": 225364072, "value": "0x3e7e6bde"}, {"index": 44593546, "value": "0x382a3aca"}, {"index": 260713159, "value": "0x567e7f7f"}, {"index": 168250303, "value": "0x25a0f084"}, {"index": 52384140, "value": "0xbfeef128"}, {"index": 223401610, "value": "0xe338abfb"}, {"index": 45554030, "value": "0x7c3b5280"}, {"index": 95410555, "value": "0x909bc5f1"}, {"index": 175039924, "value": "0xd8b74b9c"}, {"index": 79171087, "value": "0x8e31a22e"}, {"index": 267580473, "value": "0x26b5f1d8"}, {"index": 24168642, "value": "0x79122c00"}, {"index": 37981670, "value": "0xcafc3340"}, {"index": 171551130, "value": "0xd5e02ea3"}, {"index": 195559979, "value": "0x1aee1afd"}, {"index": 204611762, "value": "0xdb090d9a"}, {"index": 140997658, "value": "0xb049f435"}, {"index": 138925853, "value": "0x4954d8ba"}, {"index": 86637313, "value": "0x03797ba0"}, {"index": 20736778, "value": "0x196eefbd"}, {"index": 219665210, "value": "0xd153412a"}, {"index": 160430336, "value": "0xbe5d2c4b"}, {"index": 264654675, "value": "0xdaa14f0e"}, {"index": 8013395, "value": "0x8e61ed07"}, {"index": 228945585, "value": "0x9e9a64c6"}, {"index": 213884386, "value": "0x2e29ff36"}, {"index": 104419827, "value": "0x392a8589"}, {"index": 44185464, "value": "0xb56a5912"}, {"index": 142737231, "value": "0xfa6e8b57"}, {"index": 99284897, "value": "0xd1a737cb"}, {"index": 132475900, "value": "0xb0fa841a"}, {"index": 61861762, "value": "0xbe1c341f"}, {"index": 132056166, "value": "0xe25be0f1"}, {"index": 262388043, "value": "0xe937f543"}, {"index": 91878046, "value": "0xebab2248"}, {"index": 117353561, "value": "0x8e1b607a"}, {"index": 124768597, "value": "0x202a2fed"}, {"index": 71352993, "value": "0x95e2819c"}, {"index": 190698941, "value": "0x9c9652d4"}, {"index": 46055428, "value": "0x32fedef0"}, {"index": 55281366, "value": "0xdecfff82"}, {"index": 165145231, "value": "0xcb5d43e5"}, {"index": 106810753, "value": "0xb735806a"}, {"index": 171985651, "value": "0x8905939c"}, {"index": 232085256, "value": "0xfbf8472d"}, {"index": 159510492, "value": "0xada74e5d"}, {"index": 40072060, "value": "0x7ebdeeea"}, {"index": 209107596, "value": "0x0119f2b3"}, {"index": 39023794, "value": "0xa9a376b8"}], + "cache_head": ["0x355a86d2", "0x7957db1c", "0xd21772af", "0x6fc1e09b", "0xd55ce61d", "0x6e6a278b", "0xd3f543ce", "0x223d8e82", "0x143ab337", "0x2e9f05bd", "0x2eb389bf", "0x0c6e449e", "0x5cfa4222", "0xba6560fe", "0x8e3e1aa4", "0xdbcc1d53"], + "cache_last_line": ["0x41190d91", "0xbd277957", "0x22ddbb49", "0x6986f207", "0xdf69a4d6", "0x26401a3a", "0x818230fb", "0xc417122d", "0x3597b211", "0xb553ce55", "0xcf39cc0d", "0x3b7fc43a", "0x3fd43b00", "0x67e1c80e", "0xffa7ea7d", "0xca2960ab"], + "cache_fnv1a64": "0x48c4f5bf24166b2e" +} diff --git a/site/index.html b/site/index.html index 16262ed13..b7743a223 100644 --- a/site/index.html +++ b/site/index.html @@ -257,7 +257,8 @@ footer .wrap{padding-block:48px 32px} mined shards being proven proven - checkpoint locked by sustained miners + checkpoint, every 30 s on chain + locked by sustained miners, final
@@ -477,7 +478,7 @@ footer .wrap{padding-block:48px 32px} function size(){W=c.clientWidth;H=c.clientHeight;c.width=W*dpr;c.height=H*dpr;x.setTransform(dpr,0,0,dpr,0,0);S=W<600?18:26;GAP=S*3.1;speed=GAP/1000;} size();window.addEventListener('resize',size); function laneY(l){return H*0.14+l*(H*0.72/(lanes-1));} - function spawn(px){n++;var lane=Math.floor(Math.random()*lanes);var b={id:n,x:px,y:laneY(lane),born:performance.now(),shards:[0,0,0,0],claimed:[false,false,false,false],state:'mined',cp:n%10===0,parents:[],glow:0}; + function spawn(px){n++;var lane=Math.floor(Math.random()*lanes);var b={id:n,x:px,y:laneY(lane),born:performance.now(),shards:[0,0,0,0],claimed:[false,false,false,false],state:'mined',cp:n%6===0,parents:[],glow:0,flash:0}; var cands=blocks.slice(-8).filter(function(q){return q.x80)blocks.shift();cb.textContent=n;} @@ -489,11 +490,12 @@ footer .wrap{padding-block:48px 32px} blocks.forEach(function(b){b.x-=speed*dt;var age=t-b.born; if(b.state==='mined'&&age>1200)b.state='proving'; if(b.state==='proving'){for(var i=0;i<4;i++){if(!b.claimed[i]&&Math.random()<0.012){b.claimed[i]=true;var fromTop=Math.random()<0.5;sparks.push({x:Math.random()*W,y:fromTop?-6:H+6,tx:b.x,ty:b.y,b:b,i:i,p:0});}} - if(b.shards.every(function(v){return v>=1;})){b.state='proven';b.glow=1;proven++;cp.textContent=proven;if(b.cp)setTimeout(function(){b.locked=true;locked++;cl.textContent=locked;ripples.push({x:b.x,y:b.y,r:S,a:1});},1800);}} + if(b.shards.every(function(v){return v>=1;})){b.state='proven';b.glow=1;proven++;cp.textContent=proven;if(b.cp)setTimeout(function(){b.locked=true;b.flash=1;locked++;cl.textContent=locked;ripples.push({x:b.x,y:b.y,r:S,a:1});ripples.push({x:b.x,y:b.y,r:S*0.5,a:1});},1800);}} + if(b.flash>0)b.flash-=0.012; if(b.glow>0)b.glow-=0.02;}); sparks.forEach(function(sp){sp.p+=0.035;sp.tx=sp.b.x;sp.ty=sp.b.y;if(sp.p>=1){sp.done=true;sp.b.shards[sp.i]=1;}}); sparks=sparks.filter(function(sp){return !sp.done;}); - ripples.forEach(function(r){r.r+=1.6;r.a-=0.018;r.x-=speed*dt;});ripples=ripples.filter(function(r){return r.a>0;}); + ripples.forEach(function(r){r.r+=2.2;r.a-=0.012;r.x-=speed*dt;});ripples=ripples.filter(function(r){return r.a>0;}); x.clearRect(0,0,W,H); // faint lane rails x.strokeStyle='rgba(42,42,48,0.6)';x.lineWidth=1;for(var l=0;llk.x))lk=b;}); - if(lk){x.strokeStyle='rgba(242,84,27,0.5)';x.setLineDash([4,6]);x.beginPath();x.moveTo(lk.x,8);x.lineTo(lk.x,H-8);x.stroke();x.setLineDash([]);} + if(lk){if(lk.flash>0){x.fillStyle='rgba(242,84,27,'+(0.16*lk.flash)+')';x.fillRect(0,0,lk.x,H);}x.strokeStyle='rgba(242,84,27,0.7)';x.lineWidth=1.5;x.setLineDash([5,6]);x.beginPath();x.moveTo(lk.x,6);x.lineTo(lk.x,H-6);x.stroke();x.setLineDash([]);x.fillStyle='rgba(242,84,27,0.95)';x.font='600 '+Math.round(S*0.42)+'px IBM Plex Mono, monospace';x.textAlign='right';x.fillText('FINAL \u2190',lk.x-6,14);} // blocks blocks.forEach(function(b){var h=S/2; if(b.glow>0){var g=x.createRadialGradient(b.x,b.y,0,b.x,b.y,S*1.6);g.addColorStop(0,'rgba(242,84,27,'+(0.45*b.glow)+')');g.addColorStop(1,'rgba(242,84,27,0)');x.fillStyle=g;x.beginPath();x.arc(b.x,b.y,S*1.6,0,Math.PI*2);x.fill();} if(b.state==='proven'){x.fillStyle='#F2541B';rr(b.x-h,b.y-h,S,S,S*0.24);x.fill();} else{x.fillStyle='#0C0C0E';rr(b.x-h,b.y-h,S,S,S*0.24);x.fill();x.strokeStyle=b.state==='proving'?'rgba(255,179,92,0.9)':'rgba(58,58,66,1)';x.lineWidth=2;rr(b.x-h,b.y-h,S,S,S*0.24);x.stroke(); if(b.state==='proving'){var q=S/2-3;for(var i=0;i<4;i++){if(b.shards[i]>=1){x.fillStyle='#FFB35C';var qx=b.x-h+3+(i%2)*q,qy=b.y-h+3+Math.floor(i/2)*q;rr(qx,qy,q-1,q-1,2);x.fill();}}}} - if(b.locked){x.strokeStyle='rgba(242,84,27,0.95)';x.lineWidth=2;x.beginPath();x.arc(b.x,b.y,S*0.95,0,Math.PI*2);x.stroke();x.fillStyle='#F4F1EC';x.font='bold '+Math.round(S*0.45)+'px IBM Plex Mono, monospace';x.textAlign='center';x.fillText('L',b.x,b.y+S*0.16);} + if(b.cp&&!b.locked){x.strokeStyle='rgba(255,179,92,0.7)';x.lineWidth=1.5;x.setLineDash([3,3]);x.beginPath();x.arc(b.x,b.y,S*1.05,0,Math.PI*2);x.stroke();x.setLineDash([]);x.fillStyle='rgba(255,179,92,0.9)';x.font='600 '+Math.round(S*0.38)+'px IBM Plex Mono, monospace';x.textAlign='center';x.fillText('CHECKPOINT',b.x,b.y-S*1.35);} + if(b.locked){var pulse=0.85+0.15*Math.sin(t/300);x.strokeStyle='rgba(242,84,27,'+pulse+')';x.lineWidth=3;x.beginPath();x.arc(b.x,b.y,S*1.05,0,Math.PI*2);x.stroke();x.fillStyle='#F2541B';rr(b.x-S*0.9,b.y-S*1.75,S*1.8,S*0.6,4);x.fill();x.fillStyle='#0C0C0E';x.font='700 '+Math.round(S*0.4)+'px IBM Plex Mono, monospace';x.textAlign='center';x.fillText('LOCKED',b.x,b.y-S*1.32);} }); // sparks (miners claiming shards) sparks.forEach(function(sp){var e=1-Math.pow(1-sp.p,3);var sx=sp.x+(sp.tx-sp.x)*e,sy=sp.y+(sp.ty-sp.y)*e;x.strokeStyle='rgba(255,179,92,'+(0.5*(1-sp.p))+')';x.lineWidth=1;x.beginPath();x.moveTo(sp.x+(sp.tx-sp.x)*Math.max(0,e-0.15),sp.y+(sp.ty-sp.y)*Math.max(0,e-0.15));x.lineTo(sx,sy);x.stroke();x.fillStyle='#FFB35C';x.beginPath();x.arc(sx,sy,2.2,0,Math.PI*2);x.fill();}); ripples.forEach(function(r){x.strokeStyle='rgba(242,84,27,'+r.a+')';x.lineWidth=1.5;x.beginPath();x.arc(r.x,r.y,r.r,0,Math.PI*2);x.stroke();}); if(!reduce)requestAnimationFrame(frame); } - for(var i=0;i5){b.state='proven';b.shards=[1,1,1,1];b.claimed=[true,true,true,true];proven++;if(b.cp&&i>8){b.locked=true;locked++;}}else if(i>1){b.state='proving';b.claimed=[true,true,false,false];b.shards=[1,1,0,0];}} + for(var i=0;i5){b.state='proven';b.shards=[1,1,1,1];b.claimed=[true,true,true,true];proven++;if(b.cp&&i>5){b.locked=true;locked++;}}else if(i>1){b.state='proving';b.claimed=[true,true,false,false];b.shards=[1,1,0,0];}} cb.textContent=n;cp.textContent=proven;cl.textContent=locked; requestAnimationFrame(frame); })();