145 lines
7.5 KiB
Text
145 lines
7.5 KiB
Text
// Generated by proto-metal/igneum-bench --export-pack for seed "igneum-hourly". 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"
|
|
|
|
__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;
|
|
}
|
|
|
|
// dataset[i] = ds_elem(i, d0, d1) for i < n. Same closed form as the Metal igneum_fill kernel.
|
|
__global__ void igneum_fill(uint32_t* ds, uint32_t n, uint32_t d0, uint32_t d1) {
|
|
uint32_t i = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (i < n) ds[i] = ds_elem(i, d0, d1);
|
|
}
|
|
|
|
// 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 ^ 0x6bdee811u; x += 0x9e3779b9u; x = splitmix32(x); r0 = x ^ 0x8f488bbeu; } // SEEDW[0], 0x9e3779b9u * 1u, SEEDW[1]
|
|
{ uint32_t x = nonce ^ 0x8f488bbeu; x += 0x3c6ef372u; x = splitmix32(x); r1 = x ^ 0xc5cdece7u; } // SEEDW[1], 0x9e3779b9u * 2u, SEEDW[2]
|
|
{ uint32_t x = nonce ^ 0xc5cdece7u; x += 0xdaa66d2bu; x = splitmix32(x); r2 = x ^ 0x210af22du; } // SEEDW[2], 0x9e3779b9u * 3u, SEEDW[3]
|
|
{ uint32_t x = nonce ^ 0x210af22du; x += 0x78dde6e4u; x = splitmix32(x); r3 = x ^ 0x2f687b65u; } // SEEDW[3], 0x9e3779b9u * 4u, SEEDW[4]
|
|
{ uint32_t x = nonce ^ 0x2f687b65u; x += 0x1715609du; x = splitmix32(x); r4 = x ^ 0x17471eeeu; } // SEEDW[4], 0x9e3779b9u * 5u, SEEDW[5]
|
|
{ uint32_t x = nonce ^ 0x17471eeeu; x += 0xb54cda56u; x = splitmix32(x); r5 = x ^ 0xee16e284u; } // SEEDW[5], 0x9e3779b9u * 6u, SEEDW[6]
|
|
{ uint32_t x = nonce ^ 0xee16e284u; x += 0x5384540fu; x = splitmix32(x); r6 = x ^ 0xfc9eb8f9u; } // SEEDW[6], 0x9e3779b9u * 7u, SEEDW[7]
|
|
{ uint32_t x = nonce ^ 0xfc9eb8f9u; x += 0xf1bbcdc8u; x = splitmix32(x); r7 = x ^ 0x6bdee811u; } // SEEDW[7], 0x9e3779b9u * 8u, SEEDW[0]
|
|
|
|
for (uint32_t it = 0u; it < 8u; ++it) {
|
|
uint32_t sel = r0;
|
|
r6 = r6 ^ ds[r5 & mask]; // 0 load
|
|
r3 = r3 ^ r7; // 1 xor
|
|
r6 = __umulhi(r6, r2); // 2 mulhi
|
|
r1 = r1 + r0 + ((((sel >> 0u) & 1u) != 0u) ? 0x43f8f369u : 0x1eb46b1cu); // 3 add
|
|
r3 = r3 ^ ds[r0 & mask]; // 4 load
|
|
r5 = r5 | r7; // 5 or
|
|
r4 = r4 ^ ds[r6 & mask]; // 6 load
|
|
r4 = rotl_imm(r4, 21u); // 7 rotl
|
|
r6 = r6 ^ ds[r3 & mask]; // 8 load
|
|
r6 = r6 ^ ds[r1 & mask]; // 9 load
|
|
r0 = r0 ^ __shfl_xor_sync(0xffffffffu, r3, 8); // 10 shfl
|
|
r2 = r2 ^ r3; // 11 xor
|
|
r2 = r2 + r7 + ((((sel >> 19u) & 1u) != 0u) ? 0xc26c7c2au : 0x3a1ce85eu); // 12 add
|
|
r4 = r4 ^ ds[r3 & mask]; // 13 load
|
|
r7 = r7 ^ ds[r1 & mask]; // 14 load
|
|
r2 = __umulhi(r2, r3); // 15 mulhi
|
|
r5 = r5 ^ ds[r2 & mask]; // 16 load
|
|
r5 = r5 ^ ds[r1 & mask]; // 17 load
|
|
r4 = r4 ^ ds[r7 & mask]; // 18 load
|
|
r2 = rotr_var(r2, r1); // 19 rotr
|
|
r7 = r7 ^ ds[r0 & mask]; // 20 load
|
|
r4 = r4 ^ ds[r6 & mask]; // 21 load
|
|
r7 = rotl_imm(r7, 25u); // 22 rotl
|
|
r3 = r3 + r5 + ((((sel >> 15u) & 1u) != 0u) ? 0x98ae0055u : 0x942b819bu); // 23 add
|
|
r3 = __umulhi(r3, r4); // 24 mulhi
|
|
r6 = r6 ^ r1; // 25 xor
|
|
r1 = rotl_imm(r1, 31u); // 26 rotl
|
|
r3 = r3 ^ __shfl_xor_sync(0xffffffffu, r6, 4); // 27 shfl
|
|
r6 = r6 - r5; // 28 sub
|
|
r6 = rotr_var(r6, r3); // 29 rotr
|
|
r0 = r0 ^ __shfl_xor_sync(0xffffffffu, r4, 4); // 30 shfl
|
|
r4 = rotl_imm(r4, 30u); // 31 rotl
|
|
r2 = r2 - r1; // 32 sub
|
|
r5 = r5 | r4; // 33 or
|
|
r7 = r6 * r3 + r7; // 34 mad
|
|
r5 = r5 * r0; // 35 mul
|
|
r5 = r5 - r3; // 36 sub
|
|
r2 = r2 + r7 + ((((sel >> 5u) & 1u) != 0u) ? 0x6b5970b5u : 0x473ecfd5u); // 37 add
|
|
r2 = r2 ^ ds[r7 & mask]; // 38 load
|
|
r2 = rotr_var(r2, r6); // 39 rotr
|
|
r0 = r0 ^ r5; // 40 xor
|
|
r4 = r4 ^ __shfl_xor_sync(0xffffffffu, r5, 8); // 41 shfl
|
|
r1 = r1 * r6; // 42 mul
|
|
r0 = r4 * r1 + r0; // 43 mad
|
|
r1 = r1 + r4 + ((((sel >> 20u) & 1u) != 0u) ? 0x4a502c22u : 0x04e78f3bu); // 44 add
|
|
r6 = r2 * r7 + r6; // 45 mad
|
|
r1 = r1 ^ r0; // 46 xor
|
|
r5 = r5 ^ ds[r7 & mask]; // 47 load
|
|
r0 = r0 | r4; // 48 or
|
|
r5 = r5 ^ __shfl_xor_sync(0xffffffffu, r4, 16); // 49 shfl
|
|
r7 = r7 + r0 + ((((sel >> 17u) & 1u) != 0u) ? 0x6fabf9ceu : 0x0a3df170u); // 50 add
|
|
r6 = r6 ^ r3; // 51 xor
|
|
r1 = r1 + r2 + ((((sel >> 21u) & 1u) != 0u) ? 0xad344ca0u : 0xc99bce6fu); // 52 add
|
|
r6 = r6 * r7; // 53 mul
|
|
r3 = __umulhi(r3, r4); // 54 mulhi
|
|
r7 = r7 * r1; // 55 mul
|
|
r7 = r7 ^ r1; // 56 xor
|
|
r2 = r2 * r7; // 57 mul
|
|
r2 = r2 ^ ds[r1 & mask]; // 58 load
|
|
r7 = r4 * r5 + r7; // 59 mad
|
|
r2 = r2 ^ ds[r7 & mask]; // 60 load
|
|
r0 = r2 * r3 + r0; // 61 mad
|
|
r1 = __umulhi(r1, r5); // 62 mulhi
|
|
r7 = r7 + r3 + ((((sel >> 24u) & 1u) != 0u) ? 0x05e4fc1du : 0xc23e27c9u); // 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;
|
|
}
|
|
|
|
// Host-side launch wrappers. Declared in program.h, called from host.cu.
|
|
cudaError_t igneum_launch_fill(uint32_t* ds, uint32_t nWords, uint32_t d0, uint32_t d1) {
|
|
if (nWords == 0u) return cudaErrorInvalidValue;
|
|
uint32_t block = 256u;
|
|
uint32_t grid = (nWords + block - 1u) / block;
|
|
igneum_fill<<<grid, block>>>(ds, nWords, d0, d1);
|
|
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);
|
|
}
|