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>
123 lines
12 KiB
Text
123 lines
12 KiB
Text
// Generated by igneum-pow export (generator v2) for seed "igneum-readwidth/B/5". Do not edit by hand.
|
|
// Header-bound twin of igneum_hash in kernel.cu: the init words come from a kernel argument, not SEEDW.
|
|
// Host declarations (also in program_bound.h if present):
|
|
// struct IgneumInitWords { uint32_t w[8]; };
|
|
// cudaError_t igneum_launch_hash_bound(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask,
|
|
// IgneumInitWords iw, uint32_t nonces, uint32_t blockWarps);
|
|
// cudaError_t igneum_hash_bound_info(int* numRegs, int* blocksPerSM, uint32_t blockWarps);
|
|
#include <cuda_runtime.h>
|
|
#include <cstdint>
|
|
#include "program.h"
|
|
|
|
struct IgneumInitWords { uint32_t w[8]; };
|
|
|
|
__device__ __forceinline__ uint32_t splitmix32(uint32_t x) {
|
|
x ^= x >> 16; x *= 0x7feb352du;
|
|
x ^= x >> 15; x *= 0x846ca68bu;
|
|
x ^= x >> 16;
|
|
return x;
|
|
}
|
|
__device__ __forceinline__ uint32_t rotl_imm(uint32_t x, uint32_t n) { return (x << n) | (x >> (32u - n)); }
|
|
__device__ __forceinline__ uint32_t rotr_var(uint32_t x, uint32_t n) { n &= 31u; return (x >> n) | (x << ((32u - n) & 31u)); }
|
|
|
|
__global__ void igneum_hash_bound(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask, IgneumInitWords iw) {
|
|
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 ^ iw.w[0]; x += 0x9e3779b9u * 1u; x = splitmix32(x); r0 = x ^ iw.w[1]; }
|
|
{ uint32_t x = nonce ^ iw.w[1]; x += 0x9e3779b9u * 2u; x = splitmix32(x); r1 = x ^ iw.w[2]; }
|
|
{ uint32_t x = nonce ^ iw.w[2]; x += 0x9e3779b9u * 3u; x = splitmix32(x); r2 = x ^ iw.w[3]; }
|
|
{ uint32_t x = nonce ^ iw.w[3]; x += 0x9e3779b9u * 4u; x = splitmix32(x); r3 = x ^ iw.w[4]; }
|
|
{ uint32_t x = nonce ^ iw.w[4]; x += 0x9e3779b9u * 5u; x = splitmix32(x); r4 = x ^ iw.w[5]; }
|
|
{ uint32_t x = nonce ^ iw.w[5]; x += 0x9e3779b9u * 6u; x = splitmix32(x); r5 = x ^ iw.w[6]; }
|
|
{ uint32_t x = nonce ^ iw.w[6]; x += 0x9e3779b9u * 7u; x = splitmix32(x); r6 = x ^ iw.w[7]; }
|
|
{ uint32_t x = nonce ^ iw.w[7]; x += 0x9e3779b9u * 8u; x = splitmix32(x); r7 = x ^ iw.w[0]; }
|
|
|
|
for (uint32_t it = 0u; it < 8u; ++it) {
|
|
uint32_t sel = r0;
|
|
r3 = r3 * r1; // 0 mul
|
|
r2 = r2 ^ __shfl_xor_sync(0xffffffffu, r4, 4); // 1 shfl
|
|
r0 = r0 ^ r6; // 2 xor
|
|
{ uint32_t b_ = (r3 & 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_ = 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; 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; r4 = x_; } // 3 load
|
|
r6 = r6 ^ ds[r0 & mask]; // 4 load
|
|
r2 = r2 ^ __shfl_xor_sync(0xffffffffu, r0, 2); // 5 shfl
|
|
r7 = r7 ^ __shfl_xor_sync(0xffffffffu, r6, 4); // 6 shfl
|
|
r5 = rotl_imm(r5, 8u); // 7 rotl
|
|
r1 = r5 * r7 + r1; // 8 mad
|
|
r0 = r0 * r3; // 9 mul
|
|
r0 = r0 * r4; // 10 mul
|
|
r3 = r3 ^ __shfl_xor_sync(0xffffffffu, r0, 16); // 11 shfl
|
|
{ uint32_t b_ = (r5 & 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_; } // 12 load
|
|
r0 = r0 ^ __shfl_xor_sync(0xffffffffu, r2, 16); // 13 shfl
|
|
r4 = r4 - r6; // 14 sub
|
|
r7 = r7 | r5; // 15 or
|
|
r0 = rotr_var(r0, r6); // 16 rotr
|
|
r1 = r1 - r7; // 17 sub
|
|
r3 = r3 - r6; // 18 sub
|
|
r5 = r5 - r4; // 19 sub
|
|
r3 = r3 ^ __shfl_xor_sync(0xffffffffu, r5, 1); // 20 shfl
|
|
r1 = r1 | r2; // 21 or
|
|
r1 = r1 ^ __shfl_xor_sync(0xffffffffu, r0, 4); // 22 shfl
|
|
r5 = rotr_var(r5, r3); // 23 rotr
|
|
{ uint32_t b_ = (r2 & 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_; } // 24 load
|
|
{ uint32_t b_ = (r0 & 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_; } // 25 load
|
|
r7 = r7 | r1; // 26 or
|
|
{ uint32_t b_ = (r7 & mask) & ~3u; const uint4* l_ = (const uint4*)(ds + b_); uint4 v0_ = l_[0]; uint32_t x_ = r2 ^ 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; r2 = x_; } // 27 load
|
|
r3 = r3 ^ r4; // 28 xor
|
|
{ uint32_t b_ = (r5 & 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_; } // 29 load
|
|
r5 = r5 + r4 + ((((sel >> 15u) & 1u) != 0u) ? 0x90aaff78u : 0xd0651e01u); // 30 add
|
|
{ uint32_t b_ = (r3 & 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_ = 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; 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; r1 = x_; } // 31 load
|
|
r0 = r0 * r4; // 32 mul
|
|
r0 = r0 + r3 + ((((sel >> 2u) & 1u) != 0u) ? 0x3d45d6cfu : 0x8ac5ddc2u); // 33 add
|
|
{ uint32_t b_ = (r4 & 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_; } // 34 load
|
|
r6 = r6 ^ r2; // 35 xor
|
|
r1 = r1 ^ r6; // 36 xor
|
|
r0 = r4 * r3 + r0; // 37 mad
|
|
r7 = rotl_imm(r7, 17u); // 38 rotl
|
|
r5 = r5 + r7 + ((((sel >> 0u) & 1u) != 0u) ? 0x4b7eb232u : 0x651ff080u); // 39 add
|
|
{ uint32_t b_ = (r2 & 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_; } // 40 load
|
|
r6 = __umulhi(r6, r1); // 41 mulhi
|
|
{ uint32_t b_ = (r6 & 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_ = 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; 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; r1 = x_; } // 42 load
|
|
{ uint32_t b_ = (r0 & 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_; } // 43 load
|
|
{ uint32_t b_ = (r7 & 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_; } // 44 load
|
|
r3 = r3 * r2; // 45 mul
|
|
r6 = r6 ^ __shfl_xor_sync(0xffffffffu, r5, 2); // 46 shfl
|
|
r4 = r2 * r3 + r4; // 47 mad
|
|
r4 = r4 + r2 + ((((sel >> 20u) & 1u) != 0u) ? 0x97f34047u : 0x4f0330b4u); // 48 add
|
|
r6 = r6 ^ r3; // 49 xor
|
|
r0 = r0 + r2 + ((((sel >> 17u) & 1u) != 0u) ? 0xb24ddcb5u : 0x6e73a5d8u); // 50 add
|
|
r6 = rotl_imm(r6, 22u); // 51 rotl
|
|
r4 = r4 * r0; // 52 mul
|
|
r6 = rotr_var(r6, r3); // 53 rotr
|
|
{ uint32_t b_ = (r6 & 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_; } // 54 load
|
|
r3 = r3 - r1; // 55 sub
|
|
r3 = r3 ^ __shfl_xor_sync(0xffffffffu, r7, 1); // 56 shfl
|
|
r3 = r3 | r1; // 57 or
|
|
r4 = r7 * r4 + r4; // 58 mad
|
|
{ uint32_t b_ = (r0 & 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_; } // 59 load
|
|
r4 = r6 * r3 + r4; // 60 mad
|
|
r2 = rotl_imm(r2, 30u); // 61 rotl
|
|
{ uint32_t b_ = (r1 & mask) & ~3u; const uint4* l_ = (const uint4*)(ds + b_); uint4 v0_ = l_[0]; uint32_t x_ = r2 ^ 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; r2 = x_; } // 62 load
|
|
r7 = r7 + r4 + ((((sel >> 3u) & 1u) != 0u) ? 0x4bcca5cau : 0x12a17f52u); // 63 add
|
|
}
|
|
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;
|
|
}
|
|
|
|
cudaError_t igneum_launch_hash_bound(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask,
|
|
IgneumInitWords iw, 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_bound<<<nonces / block, block>>>(ds, out, baseNonce, mask, iw);
|
|
return cudaGetLastError();
|
|
}
|
|
|
|
cudaError_t igneum_hash_bound_info(int* numRegs, int* blocksPerSM, uint32_t blockWarps) {
|
|
cudaFuncAttributes attr;
|
|
cudaError_t e = cudaFuncGetAttributes(&attr, igneum_hash_bound);
|
|
if (e != cudaSuccess) return e;
|
|
*numRegs = attr.numRegs;
|
|
return cudaOccupancyMaxActiveBlocksPerMultiprocessor(blocksPerSM, igneum_hash_bound, (int)(32u * blockWarps), 0);
|
|
}
|