igneum/proto-newpow/state-dataset/kernel.cu

// Generated by igneum-pow export (generator v2) 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;
        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;
}

// 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);
}