Chain scene: checkpoints tagged from birth, locked tag, flash wash and final label
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
This commit is contained in:
parent
1321907aa6
commit
c4e63c4419
9 changed files with 727 additions and 7 deletions
162
proto-cuda/packs/igneum-genesis-mh/kernel.cu
Normal file
162
proto-cuda/packs/igneum-genesis-mh/kernel.cu
Normal file
|
|
@ -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 <cuda_runtime.h>
|
||||
#include <cstdint>
|
||||
#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<<<grid, block>>>(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<<<grid, block>>>(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<<<nonces / block, block>>>(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);
|
||||
}
|
||||
101
proto-cuda/packs/igneum-genesis-mh/memhard.h
Normal file
101
proto-cuda/packs/igneum-genesis-mh/memhard.h
Normal file
|
|
@ -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 <cstdint>
|
||||
#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]; }
|
||||
106
proto-cuda/packs/igneum-genesis-mh/memhard.metal
Normal file
106
proto-cuda/packs/igneum-genesis-mh/memhard.metal
Normal file
|
|
@ -0,0 +1,106 @@
|
|||
#include <metal_stdlib>
|
||||
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];
|
||||
}
|
||||
37
proto-cuda/packs/igneum-genesis-mh/program.h
Normal file
37
proto-cuda/packs/igneum-genesis-mh/program.h
Normal file
|
|
@ -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 <cuda_runtime.h>
|
||||
#include <cstdint>
|
||||
|
||||
#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);
|
||||
114
proto-cuda/packs/igneum-genesis-mh/program.json
Normal file
114
proto-cuda/packs/igneum-genesis-mh/program.json
Normal file
|
|
@ -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}
|
||||
]
|
||||
}
|
||||
108
proto-cuda/packs/igneum-genesis-mh/program.metal
Normal file
108
proto-cuda/packs/igneum-genesis-mh/program.metal
Normal file
|
|
@ -0,0 +1,108 @@
|
|||
#include <metal_stdlib>
|
||||
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;
|
||||
}
|
||||
53
proto-cuda/packs/igneum-genesis-mh/vectors.h
Normal file
53
proto-cuda/packs/igneum-genesis-mh/vectors.h
Normal file
|
|
@ -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 <cstdint>
|
||||
|
||||
#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;
|
||||
36
proto-cuda/packs/igneum-genesis-mh/vectors.json
Normal file
36
proto-cuda/packs/igneum-genesis-mh/vectors.json
Normal file
|
|
@ -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"
|
||||
}
|
||||
|
|
@ -257,7 +257,8 @@ footer .wrap{padding-block:48px 32px}
|
|||
<span><i class="sw" style="border:2px solid var(--line-2)"></i>mined</span>
|
||||
<span><i class="sw" style="background:var(--molten)"></i>shards being proven</span>
|
||||
<span><i class="sw" style="background:var(--ember)"></i>proven</span>
|
||||
<span><i class="sw" style="border:2px solid var(--ember);border-radius:50%"></i>checkpoint locked by sustained miners</span>
|
||||
<span><i class="sw" style="border:2px dashed var(--molten);border-radius:50%"></i>checkpoint, every 30 s on chain</span>
|
||||
<span><i class="sw" style="border:3px solid var(--ember);border-radius:50%"></i>locked by sustained miners, final</span>
|
||||
</div>
|
||||
</div>
|
||||
<div class="grid" style="margin-top:24px">
|
||||
|
|
@ -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.x<px-S;}).sort(function(a2,b2){return b2.x-a2.x;});
|
||||
if(cands.length)b.parents.push(cands[0]);cands.slice(1).forEach(function(q){if(Math.random()<0.4&&b.parents.length<3)b.parents.push(q);});
|
||||
blocks.push(b);if(blocks.length>80)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;l<lanes;l++){x.beginPath();x.moveTo(0,laneY(l));x.lineTo(W,laneY(l));x.stroke();}
|
||||
|
|
@ -501,21 +503,22 @@ footer .wrap{padding-block:48px 32px}
|
|||
blocks.forEach(function(b){b.parents.forEach(function(q){if(blocks.indexOf(q)<0)return;var both=b.state==='proven'&&q.state==='proven';x.strokeStyle=both?'rgba(242,84,27,0.45)':'rgba(154,154,158,0.22)';x.lineWidth=both?1.5:1;x.beginPath();x.moveTo(b.x,b.y);var mx=(b.x+q.x)/2;x.bezierCurveTo(mx,b.y,mx,q.y,q.x,q.y);x.stroke();});});
|
||||
// lock line: everything left of the newest locked block is final
|
||||
var lk=null;blocks.forEach(function(b){if(b.locked&&(!lk||b.x>lk.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;i<Math.ceil(W/(S*3.1))+1;i++){spawn(W-i*S*3.1);var b=blocks[blocks.length-1];b.born=performance.now()-i*1000;if(i>5){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;i<Math.ceil(W/(S*3.1))+1;i++){spawn(W-i*S*3.1);var b=blocks[blocks.length-1];b.born=performance.now()-i*1000;if(i>5){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);
|
||||
})();
|
||||
|
|
|
|||
Loading…
Reference in a new issue