dc84789e34
Nothing changes for the default class: the pinned packs are byte-identical (tests/packs.rs), the v2 draw stream is untouched.
LoadClass {mix, load_slots, scratch}: fixed widths w16, w64, w64x4 (4 loads of 64 B), era mixes 50/35/15 and 25/50/25 drawn per load with one extra below(100) roll, and the scratch variant scr0/2/4/8 (persistent warps, 1 MiB per warp, tagged lazy fill, measurement only). A wide load reads the W-aligned address and folds every word: x = dst ^ w0; x = (rotl(x, 11) * 0x9e3779b1) ^ w[j]. Program ids carry the class. proto-opencl/host.c taken from opencl-rdna4 23810df (--memprobe, select read-back).
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
164 lines
13 KiB
Text
164 lines
13 KiB
Text
// Generated by igneum-pow export (generator v2) for seed "igneum-readwidth/B/0". 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 ^ 0x3673211cu; x += 0x9e3779b9u; x = splitmix32(x); r0 = x ^ 0xaae550b4u; } // SEEDW[0], 0x9e3779b9u * 1u, SEEDW[1]
|
|
{ uint32_t x = nonce ^ 0xaae550b4u; x += 0x3c6ef372u; x = splitmix32(x); r1 = x ^ 0x5a0ce2e0u; } // SEEDW[1], 0x9e3779b9u * 2u, SEEDW[2]
|
|
{ uint32_t x = nonce ^ 0x5a0ce2e0u; x += 0xdaa66d2bu; x = splitmix32(x); r2 = x ^ 0x1d2471cfu; } // SEEDW[2], 0x9e3779b9u * 3u, SEEDW[3]
|
|
{ uint32_t x = nonce ^ 0x1d2471cfu; x += 0x78dde6e4u; x = splitmix32(x); r3 = x ^ 0xba944366u; } // SEEDW[3], 0x9e3779b9u * 4u, SEEDW[4]
|
|
{ uint32_t x = nonce ^ 0xba944366u; x += 0x1715609du; x = splitmix32(x); r4 = x ^ 0xbdd4d1dbu; } // SEEDW[4], 0x9e3779b9u * 5u, SEEDW[5]
|
|
{ uint32_t x = nonce ^ 0xbdd4d1dbu; x += 0xb54cda56u; x = splitmix32(x); r5 = x ^ 0xcb1984a9u; } // SEEDW[5], 0x9e3779b9u * 6u, SEEDW[6]
|
|
{ uint32_t x = nonce ^ 0xcb1984a9u; x += 0x5384540fu; x = splitmix32(x); r6 = x ^ 0x081ce12au; } // SEEDW[6], 0x9e3779b9u * 7u, SEEDW[7]
|
|
{ uint32_t x = nonce ^ 0x081ce12au; x += 0xf1bbcdc8u; x = splitmix32(x); r7 = x ^ 0x3673211cu; } // SEEDW[7], 0x9e3779b9u * 8u, SEEDW[0]
|
|
|
|
for (uint32_t it = 0u; it < 8u; ++it) {
|
|
uint32_t sel = r0;
|
|
r0 = r0 ^ __shfl_xor_sync(0xffffffffu, r2, 1); // 0 shfl
|
|
r2 = r2 * r0; // 1 mul
|
|
r0 = r0 + r6 + ((((sel >> 20u) & 1u) != 0u) ? 0x03fa29f3u : 0x7fbf4ae6u); // 2 add
|
|
r2 = r2 * r7; // 3 mul
|
|
r6 = r6 ^ ds[r2 & mask]; // 4 load
|
|
{ uint32_t b_ = (r0 & mask) & ~3u; const uint4* l_ = (const uint4*)(ds + b_); uint4 v0_ = l_[0]; uint32_t x_ = r7 ^ v0_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.w; r7 = x_; } // 5 load
|
|
r6 = r6 ^ __shfl_xor_sync(0xffffffffu, r4, 16); // 6 shfl
|
|
r3 = r3 ^ r5; // 7 xor
|
|
r6 = r6 ^ ds[r7 & mask]; // 8 load
|
|
r0 = r0 + r7 + ((((sel >> 25u) & 1u) != 0u) ? 0x308c81bfu : 0xd59b21b0u); // 9 add
|
|
r5 = rotr_var(r5, r7); // 10 rotr
|
|
{ uint32_t b_ = (r6 & mask) & ~3u; const uint4* l_ = (const uint4*)(ds + b_); uint4 v0_ = l_[0]; uint32_t x_ = r3 ^ v0_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.w; r3 = x_; } // 11 load
|
|
r2 = r2 * r6; // 12 mul
|
|
{ uint32_t b_ = (r0 & mask) & ~15u; const uint4* l_ = (const uint4*)(ds + b_); uint4 v0_ = l_[0]; uint4 v1_ = l_[1]; uint4 v2_ = l_[2]; uint4 v3_ = l_[3]; uint32_t x_ = r7 ^ v0_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.w; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.w; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.w; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.w; r7 = x_; } // 13 load
|
|
r2 = r5 * r2 + r2; // 14 mad
|
|
r4 = r4 + r0 + ((((sel >> 24u) & 1u) != 0u) ? 0x67081e7eu : 0x902dc661u); // 15 add
|
|
{ uint32_t b_ = (r4 & mask) & ~3u; const uint4* l_ = (const uint4*)(ds + b_); uint4 v0_ = l_[0]; uint32_t x_ = r3 ^ v0_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.w; r3 = x_; } // 16 load
|
|
r3 = __umulhi(r3, r4); // 17 mulhi
|
|
{ uint32_t b_ = (r7 & mask) & ~3u; const uint4* l_ = (const uint4*)(ds + b_); uint4 v0_ = l_[0]; uint32_t x_ = r4 ^ v0_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.w; r4 = x_; } // 18 load
|
|
r2 = r2 + r1 + ((((sel >> 9u) & 1u) != 0u) ? 0xd1e74db0u : 0x3ee3182cu); // 19 add
|
|
r7 = r7 ^ r6; // 20 xor
|
|
r5 = r5 ^ r3; // 21 xor
|
|
r6 = r6 ^ __shfl_xor_sync(0xffffffffu, r2, 16); // 22 shfl
|
|
{ uint32_t b_ = (r2 & mask) & ~3u; const uint4* l_ = (const uint4*)(ds + b_); uint4 v0_ = l_[0]; uint32_t x_ = r0 ^ v0_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.w; r0 = x_; } // 23 load
|
|
r2 = r2 + r5 + ((((sel >> 22u) & 1u) != 0u) ? 0x62ac9e52u : 0xd2d451c6u); // 24 add
|
|
{ uint32_t b_ = (r7 & mask) & ~3u; const uint4* l_ = (const uint4*)(ds + b_); uint4 v0_ = l_[0]; uint32_t x_ = r3 ^ v0_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.w; r3 = x_; } // 25 load
|
|
r4 = r4 ^ __shfl_xor_sync(0xffffffffu, r1, 16); // 26 shfl
|
|
{ uint32_t b_ = (r4 & mask) & ~3u; const uint4* l_ = (const uint4*)(ds + b_); uint4 v0_ = l_[0]; uint32_t x_ = r1 ^ v0_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.w; r1 = x_; } // 27 load
|
|
r1 = r1 + r5 + ((((sel >> 2u) & 1u) != 0u) ? 0x072cfabeu : 0x111ff813u); // 28 add
|
|
r4 = r1 * r2 + r4; // 29 mad
|
|
r2 = r2 * r0; // 30 mul
|
|
r0 = r0 ^ r5; // 31 xor
|
|
r1 = r1 ^ ds[r0 & mask]; // 32 load
|
|
r2 = r2 - r3; // 33 sub
|
|
r2 = r2 ^ __shfl_xor_sync(0xffffffffu, r0, 2); // 34 shfl
|
|
r0 = r0 - r6; // 35 sub
|
|
r4 = r4 * r7; // 36 mul
|
|
r5 = r5 ^ r6; // 37 xor
|
|
r0 = __umulhi(r0, r6); // 38 mulhi
|
|
r5 = r5 ^ __shfl_xor_sync(0xffffffffu, r7, 4); // 39 shfl
|
|
{ uint32_t b_ = (r3 & mask) & ~3u; const uint4* l_ = (const uint4*)(ds + b_); uint4 v0_ = l_[0]; uint32_t x_ = r6 ^ v0_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.w; r6 = x_; } // 40 load
|
|
r2 = r2 ^ ds[r4 & mask]; // 41 load
|
|
r5 = r5 ^ ds[r6 & mask]; // 42 load
|
|
r1 = rotr_var(r1, r4); // 43 rotr
|
|
r0 = r0 | r5; // 44 or
|
|
r0 = r0 ^ __shfl_xor_sync(0xffffffffu, r2, 2); // 45 shfl
|
|
r7 = r7 * r4; // 46 mul
|
|
r3 = r3 ^ r4; // 47 xor
|
|
r2 = __umulhi(r2, r6); // 48 mulhi
|
|
r5 = r5 ^ r4; // 49 xor
|
|
r1 = rotr_var(r1, r6); // 50 rotr
|
|
r7 = r7 + r0 + ((((sel >> 22u) & 1u) != 0u) ? 0x0fe79cecu : 0x68d3a5c0u); // 51 add
|
|
r5 = r5 + r4 + ((((sel >> 8u) & 1u) != 0u) ? 0x04309e6fu : 0xeb764d91u); // 52 add
|
|
r7 = r7 ^ r0; // 53 xor
|
|
r4 = r4 + r0 + ((((sel >> 6u) & 1u) != 0u) ? 0xba0b81c6u : 0x1390b188u); // 54 add
|
|
r0 = r0 + r5 + ((((sel >> 15u) & 1u) != 0u) ? 0x5ec96dd8u : 0x4c21a6cdu); // 55 add
|
|
r7 = r5 * r3 + r7; // 56 mad
|
|
{ uint32_t b_ = (r0 & mask) & ~15u; const uint4* l_ = (const uint4*)(ds + b_); uint4 v0_ = l_[0]; uint4 v1_ = l_[1]; uint4 v2_ = l_[2]; uint4 v3_ = l_[3]; uint32_t x_ = r6 ^ v0_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.w; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.w; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.w; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.w; r6 = x_; } // 57 load
|
|
r5 = r5 + r1 + ((((sel >> 1u) & 1u) != 0u) ? 0x6c0ac4ddu : 0x68297b06u); // 58 add
|
|
r6 = rotr_var(r6, r7); // 59 rotr
|
|
{ uint32_t b_ = (r4 & mask) & ~15u; const uint4* l_ = (const uint4*)(ds + b_); uint4 v0_ = l_[0]; uint4 v1_ = l_[1]; uint4 v2_ = l_[2]; uint4 v3_ = l_[3]; uint32_t x_ = r7 ^ v0_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.w; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.w; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.w; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.w; r7 = x_; } // 60 load
|
|
r6 = r6 ^ __shfl_xor_sync(0xffffffffu, r3, 4); // 61 shfl
|
|
r4 = rotl_imm(r4, 6u); // 62 rotl
|
|
r1 = r2 * r1 + r1; // 63 mad
|
|
}
|
|
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);
|
|
}
|