50ebdd64d9
A prototype behind a new LoadClass field (derive_len) and Shape field, Shape::for_class_day: every mixer slot of the item derivation runs a straight-line program of 736 instructions drawn from the day key stream (the same SplitMix64 stream, after the 40 mixer draws), twelve two-register forms, the chain rule of SuperscalarHash made strict (every instruction reads the register the previous one wrote), an acceptance test with the x8 mixer's operation and multiply counts from the code as floors (72 x 144 as written, 72 x 128 hoisted, 1,152 multiplies). The verifier runs the program with a word-major (SoA) interpreter over the 32 items of a load, dispatching on instruction pairs; no JIT. The emitter writes mh_round_0..8 into memhard.h, memhard.metal and kernel.cl. Packs dr736-genesis and dr736-devnet-epoch0 under proto-cuda/packs-ca3-derive. The v2 and v3 paths are untouched: every pinned pack re-exports byte for byte (tests/packs.rs), cargo test -p igneum-pow 58 + 7 + 4 + 19 + 7 green. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
123 lines
6.7 KiB
Text
123 lines
6.7 KiB
Text
// Generated by igneum-pow export (generator v2) for seed "igneum-epoch/edc4fa844da9dc98d37e965176f6558a31560e40502ab3ae5491b21aaaabfb07/day/69676e65756d2d6461792ffa50000000000000". 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;
|
|
r4 = r4 + r5 + ((((sel >> 13u) & 1u) != 0u) ? 0x5810667au : 0xea86e152u); // 0 add
|
|
r2 = r2 ^ __shfl_xor_sync(0xffffffffu, r0, 4); // 1 shfl
|
|
r3 = r3 + r2 + ((((sel >> 10u) & 1u) != 0u) ? 0x642e66dbu : 0x2cccb6cau); // 2 add
|
|
r0 = rotl_imm(r0, 19u); // 3 rotl
|
|
r7 = rotr_var(r7, r6); // 4 rotr
|
|
r7 = r7 + r4 + ((((sel >> 21u) & 1u) != 0u) ? 0xc1535555u : 0xee02465fu); // 5 add
|
|
r1 = __umulhi(r1, r7); // 6 mulhi
|
|
r4 = r4 ^ ds[r2 & mask]; // 7 load
|
|
r7 = r7 ^ ds[r4 & mask]; // 8 load
|
|
r0 = r0 ^ ds[r3 & mask]; // 9 load
|
|
r5 = r5 ^ ds[r1 & mask]; // 10 load
|
|
r1 = r1 ^ ds[r5 & mask]; // 11 load
|
|
r3 = __umulhi(r3, r5); // 12 mulhi
|
|
r1 = r1 ^ ds[r3 & mask]; // 13 load
|
|
r0 = r0 - r3; // 14 sub
|
|
r5 = r1 * r3 + r5; // 15 mad
|
|
r6 = __umulhi(r6, r1); // 16 mulhi
|
|
r5 = r5 + r2 + ((((sel >> 28u) & 1u) != 0u) ? 0x8b965b57u : 0x697b3d00u); // 17 add
|
|
r0 = __umulhi(r0, r6); // 18 mulhi
|
|
r5 = rotr_var(r5, r3); // 19 rotr
|
|
r5 = __umulhi(r5, r2); // 20 mulhi
|
|
r1 = r1 + r0 + ((((sel >> 1u) & 1u) != 0u) ? 0x6d7e8d05u : 0xebcf247au); // 21 add
|
|
r7 = r7 + r5 + ((((sel >> 12u) & 1u) != 0u) ? 0xb9e3577eu : 0xf66e7017u); // 22 add
|
|
r1 = __umulhi(r1, r5); // 23 mulhi
|
|
r2 = r2 - r5; // 24 sub
|
|
r7 = r7 + r4 + ((((sel >> 2u) & 1u) != 0u) ? 0x699ef1bbu : 0x08ffa6c7u); // 25 add
|
|
r3 = r3 ^ __shfl_xor_sync(0xffffffffu, r4, 2); // 26 shfl
|
|
r7 = r7 + r1 + ((((sel >> 14u) & 1u) != 0u) ? 0xb4ead2fbu : 0xe60fea84u); // 27 add
|
|
r3 = r3 + r1 + ((((sel >> 6u) & 1u) != 0u) ? 0x8f30d21du : 0x65c76dabu); // 28 add
|
|
r2 = r2 ^ ds[r1 & mask]; // 29 load
|
|
r5 = r5 ^ ds[r7 & mask]; // 30 load
|
|
r2 = r2 ^ ds[r5 & mask]; // 31 load
|
|
r1 = r1 ^ __shfl_xor_sync(0xffffffffu, r7, 4); // 32 shfl
|
|
r4 = r5 * r7 + r4; // 33 mad
|
|
r4 = r4 + r2 + ((((sel >> 21u) & 1u) != 0u) ? 0xc7ce690cu : 0x0480debeu); // 34 add
|
|
r3 = r3 ^ __shfl_xor_sync(0xffffffffu, r7, 8); // 35 shfl
|
|
r7 = r7 + r1 + ((((sel >> 2u) & 1u) != 0u) ? 0xe10c2c95u : 0xc53b542eu); // 36 add
|
|
r5 = r5 ^ r7; // 37 xor
|
|
r2 = r2 | r1; // 38 or
|
|
r1 = __umulhi(r1, r0); // 39 mulhi
|
|
r6 = rotl_imm(r6, 19u); // 40 rotl
|
|
r4 = __umulhi(r4, r6); // 41 mulhi
|
|
r6 = r6 - r0; // 42 sub
|
|
r6 = r6 ^ __shfl_xor_sync(0xffffffffu, r3, 4); // 43 shfl
|
|
r4 = r4 ^ ds[r2 & mask]; // 44 load
|
|
r1 = r1 ^ r3; // 45 xor
|
|
r7 = r7 ^ ds[r0 & mask]; // 46 load
|
|
r3 = r3 ^ ds[r1 & mask]; // 47 load
|
|
r5 = r5 * r3; // 48 mul
|
|
r1 = r1 - r5; // 49 sub
|
|
r2 = rotl_imm(r2, 8u); // 50 rotl
|
|
r1 = r1 + r5 + ((((sel >> 23u) & 1u) != 0u) ? 0x77b9bd43u : 0xa900fec4u); // 51 add
|
|
r4 = r4 ^ ds[r7 & mask]; // 52 load
|
|
r2 = r2 - r7; // 53 sub
|
|
r4 = r4 ^ r0; // 54 xor
|
|
r1 = r1 + r6 + ((((sel >> 14u) & 1u) != 0u) ? 0x83e825bfu : 0xe09f54e9u); // 55 add
|
|
r2 = r2 ^ ds[r4 & mask]; // 56 load
|
|
r0 = r1 * r4 + r0; // 57 mad
|
|
r3 = r3 ^ ds[r5 & mask]; // 58 load
|
|
r5 = r5 | r6; // 59 or
|
|
r6 = r5 * r7 + r6; // 60 mad
|
|
r4 = rotl_imm(r4, 28u); // 61 rotl
|
|
r5 = __umulhi(r5, r0); // 62 mulhi
|
|
r3 = r3 ^ ds[r6 & mask]; // 63 load
|
|
}
|
|
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);
|
|
}
|