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
6.4 KiB
Text
123 lines
6.4 KiB
Text
// Generated by igneum-pow export (generator v2) for seed "igneum-genesis". 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;
|
|
r2 = r3 * r4 + r2; // 0 mad
|
|
r2 = r1 * r1 + r2; // 1 mad
|
|
r2 = r3 * r2 + r2; // 2 mad
|
|
r3 = r3 ^ r5; // 3 xor
|
|
r7 = r7 ^ ds[r2 & mask]; // 4 load
|
|
r5 = r5 ^ ds[r7 & mask]; // 5 load
|
|
r1 = r1 ^ __shfl_xor_sync(0xffffffffu, r4, 8); // 6 shfl
|
|
r7 = r7 ^ __shfl_xor_sync(0xffffffffu, r3, 8); // 7 shfl
|
|
r1 = __umulhi(r1, r5); // 8 mulhi
|
|
r6 = rotr_var(r6, r3); // 9 rotr
|
|
r3 = r3 | r4; // 10 or
|
|
r4 = r4 ^ ds[r3 & mask]; // 11 load
|
|
r0 = __umulhi(r0, r4); // 12 mulhi
|
|
r5 = r5 + r1 + ((((sel >> 30u) & 1u) != 0u) ? 0xd3177981u : 0xc7934706u); // 13 add
|
|
r0 = r0 ^ ds[r4 & mask]; // 14 load
|
|
r2 = r2 - r4; // 15 sub
|
|
r2 = r2 ^ ds[r0 & mask]; // 16 load
|
|
r7 = r7 ^ ds[r2 & mask]; // 17 load
|
|
r7 = r7 ^ __shfl_xor_sync(0xffffffffu, r3, 4); // 18 shfl
|
|
r5 = r5 * r0; // 19 mul
|
|
r3 = r3 ^ __shfl_xor_sync(0xffffffffu, r4, 2); // 20 shfl
|
|
r2 = r2 ^ __shfl_xor_sync(0xffffffffu, r4, 16); // 21 shfl
|
|
r6 = __umulhi(r6, r2); // 22 mulhi
|
|
r6 = r6 ^ ds[r1 & mask]; // 23 load
|
|
r5 = r5 * r0; // 24 mul
|
|
r5 = rotl_imm(r5, 19u); // 25 rotl
|
|
r7 = r7 ^ __shfl_xor_sync(0xffffffffu, r6, 2); // 26 shfl
|
|
r0 = r0 ^ r5; // 27 xor
|
|
r0 = r0 ^ r4; // 28 xor
|
|
r3 = r3 - r0; // 29 sub
|
|
r5 = r5 * r1; // 30 mul
|
|
r7 = r7 ^ ds[r2 & mask]; // 31 load
|
|
r1 = r1 ^ ds[r0 & mask]; // 32 load
|
|
r5 = r5 ^ r6; // 33 xor
|
|
r5 = r5 ^ ds[r1 & mask]; // 34 load
|
|
r0 = __umulhi(r0, r5); // 35 mulhi
|
|
r5 = r5 ^ __shfl_xor_sync(0xffffffffu, r2, 4); // 36 shfl
|
|
r7 = r7 ^ ds[r0 & mask]; // 37 load
|
|
r3 = r3 + r1 + ((((sel >> 27u) & 1u) != 0u) ? 0x230c005cu : 0x75ba2fadu); // 38 add
|
|
r1 = r1 ^ __shfl_xor_sync(0xffffffffu, r5, 4); // 39 shfl
|
|
r2 = r2 ^ r5; // 40 xor
|
|
r3 = r6 * r3 + r3; // 41 mad
|
|
r6 = r6 - r7; // 42 sub
|
|
r7 = r7 ^ r0; // 43 xor
|
|
r1 = r1 ^ ds[r7 & mask]; // 44 load
|
|
r2 = r2 * r3; // 45 mul
|
|
r1 = __umulhi(r1, r5); // 46 mulhi
|
|
r4 = r4 - r3; // 47 sub
|
|
r2 = rotr_var(r2, r6); // 48 rotr
|
|
r3 = r3 ^ ds[r5 & mask]; // 49 load
|
|
r1 = r1 + r5 + ((((sel >> 7u) & 1u) != 0u) ? 0x1907970cu : 0x81b8bc2cu); // 50 add
|
|
r0 = r0 * r2; // 51 mul
|
|
r0 = r0 + r2 + ((((sel >> 6u) & 1u) != 0u) ? 0x699fd448u : 0x4f92b968u); // 52 add
|
|
r1 = r1 + r0 + ((((sel >> 12u) & 1u) != 0u) ? 0x77b1520du : 0x2bb965afu); // 53 add
|
|
r7 = rotl_imm(r7, 14u); // 54 rotl
|
|
r3 = r3 + r7 + ((((sel >> 1u) & 1u) != 0u) ? 0xa54c55a0u : 0x7b0fe07au); // 55 add
|
|
r6 = r6 ^ ds[r7 & mask]; // 56 load
|
|
r1 = rotr_var(r1, r5); // 57 rotr
|
|
r5 = r5 ^ ds[r4 & mask]; // 58 load
|
|
r6 = r6 ^ ds[r2 & mask]; // 59 load
|
|
r3 = r5 * r0 + r3; // 60 mad
|
|
r5 = r5 + r7 + ((((sel >> 31u) & 1u) != 0u) ? 0xad7493e7u : 0xaf9dd72du); // 61 add
|
|
r4 = r4 + r6 + ((((sel >> 27u) & 1u) != 0u) ? 0x1e07c3d9u : 0x89841d87u); // 62 add
|
|
r5 = rotl_imm(r5, 19u); // 63 rotl
|
|
}
|
|
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);
|
|
}
|