51dba1ce76
A shadow block of S ALU instructions run R times at the end of every iteration, drawn from the program stream after the 64 base instructions, behind LoadClass::shadow: v2 and v3 draw nothing and emit nothing (the pinned packs are byte-identical, cargo test 54 + 4 + 19 + 7 green). The interpreter, the three kernel dialects (both kernels each), program.h and program.json carry it; the acceptance rule interprets the base program only. Packs for seed igneum-genesis over class mx8 at 4,096 to 180,224 shadow instructions per hash (proto-cuda/packs-ca3-shadow), and the PC 2 bench playbook tools/ca3-shadow/pc2-shadow-bench.ps1 (passes the publisher's three checks). Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
382 lines
21 KiB
Text
382 lines
21 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
|
|
// latency-shadow block (Counter ASIC 3.0 item 8): 256 ALU instructions x 88 passes after instruction 63, no load
|
|
for (uint32_t sh = 0u; sh < 88u; ++sh) {
|
|
r5 = r5 + r2 + ((((sel >> 18u) & 1u) != 0u) ? 0x8bc12da9u : 0x92199f99u); // s0 add
|
|
r0 = r0 + r7 + ((((sel >> 26u) & 1u) != 0u) ? 0x63079e5au : 0x8d72d3adu); // s1 add
|
|
r6 = r6 ^ __shfl_xor_sync(0xffffffffu, r3, 2); // s2 shfl
|
|
r4 = r4 - r2; // s3 sub
|
|
r7 = r7 + r0 + ((((sel >> 31u) & 1u) != 0u) ? 0x5d4c7a60u : 0xb21b4babu); // s4 add
|
|
r0 = rotl_imm(r0, 11u); // s5 rotl
|
|
r0 = r0 + r1 + ((((sel >> 16u) & 1u) != 0u) ? 0xc88e2942u : 0x2fe0e98bu); // s6 add
|
|
r7 = r7 ^ r1; // s7 xor
|
|
r1 = r6 * r5 + r1; // s8 mad
|
|
r6 = r6 ^ __shfl_xor_sync(0xffffffffu, r1, 4); // s9 shfl
|
|
r1 = r2 * r2 + r1; // s10 mad
|
|
r5 = r0 * r3 + r5; // s11 mad
|
|
r2 = r2 ^ __shfl_xor_sync(0xffffffffu, r6, 4); // s12 shfl
|
|
r1 = r1 + r5 + ((((sel >> 25u) & 1u) != 0u) ? 0xbc48c63eu : 0xbdf8f9a5u); // s13 add
|
|
r1 = rotl_imm(r1, 29u); // s14 rotl
|
|
r1 = r1 - r4; // s15 sub
|
|
r7 = r7 | r1; // s16 or
|
|
r2 = r2 ^ __shfl_xor_sync(0xffffffffu, r4, 8); // s17 shfl
|
|
r7 = r7 ^ r4; // s18 xor
|
|
r6 = r6 * r1; // s19 mul
|
|
r5 = r6 * r0 + r5; // s20 mad
|
|
r3 = r3 - r1; // s21 sub
|
|
r6 = r6 * r0; // s22 mul
|
|
r2 = r2 + r0 + ((((sel >> 1u) & 1u) != 0u) ? 0x96e8f127u : 0x45c37cecu); // s23 add
|
|
r6 = r6 - r4; // s24 sub
|
|
r7 = r3 * r4 + r7; // s25 mad
|
|
r3 = rotl_imm(r3, 9u); // s26 rotl
|
|
r2 = r2 - r1; // s27 sub
|
|
r6 = __umulhi(r6, r0); // s28 mulhi
|
|
r2 = r2 ^ __shfl_xor_sync(0xffffffffu, r4, 2); // s29 shfl
|
|
r1 = r1 ^ __shfl_xor_sync(0xffffffffu, r6, 1); // s30 shfl
|
|
r1 = r1 + r6 + ((((sel >> 1u) & 1u) != 0u) ? 0xb1cdb2abu : 0x37985632u); // s31 add
|
|
r0 = rotr_var(r0, r1); // s32 rotr
|
|
r3 = r3 ^ __shfl_xor_sync(0xffffffffu, r4, 2); // s33 shfl
|
|
r0 = r0 ^ __shfl_xor_sync(0xffffffffu, r3, 4); // s34 shfl
|
|
r7 = r7 * r0; // s35 mul
|
|
r3 = r3 + r1 + ((((sel >> 0u) & 1u) != 0u) ? 0x856e0180u : 0x804e777eu); // s36 add
|
|
r1 = r1 ^ r6; // s37 xor
|
|
r2 = r2 * r7; // s38 mul
|
|
r6 = r6 + r5 + ((((sel >> 2u) & 1u) != 0u) ? 0xe9bd0cc4u : 0x0b06c8a7u); // s39 add
|
|
r5 = r5 * r7; // s40 mul
|
|
r2 = __umulhi(r2, r3); // s41 mulhi
|
|
r2 = r2 ^ r0; // s42 xor
|
|
r0 = rotl_imm(r0, 4u); // s43 rotl
|
|
r4 = __umulhi(r4, r3); // s44 mulhi
|
|
r6 = r6 + r7 + ((((sel >> 12u) & 1u) != 0u) ? 0xcdcb63f6u : 0xee822e17u); // s45 add
|
|
r3 = r2 * r0 + r3; // s46 mad
|
|
r4 = r4 ^ __shfl_xor_sync(0xffffffffu, r3, 8); // s47 shfl
|
|
r6 = __umulhi(r6, r5); // s48 mulhi
|
|
r0 = r0 - r2; // s49 sub
|
|
r3 = r3 - r5; // s50 sub
|
|
r1 = rotr_var(r1, r4); // s51 rotr
|
|
r6 = r7 * r7 + r6; // s52 mad
|
|
r5 = r5 ^ r3; // s53 xor
|
|
r1 = r1 - r0; // s54 sub
|
|
r5 = r5 - r6; // s55 sub
|
|
r3 = r3 + r2 + ((((sel >> 10u) & 1u) != 0u) ? 0xd4758987u : 0x3dfad1b6u); // s56 add
|
|
r4 = r4 + r1 + ((((sel >> 0u) & 1u) != 0u) ? 0x0602d6beu : 0xfca75bc2u); // s57 add
|
|
r5 = __umulhi(r5, r6); // s58 mulhi
|
|
r2 = r2 ^ __shfl_xor_sync(0xffffffffu, r1, 2); // s59 shfl
|
|
r2 = rotl_imm(r2, 9u); // s60 rotl
|
|
r4 = r4 | r6; // s61 or
|
|
r6 = rotr_var(r6, r4); // s62 rotr
|
|
r2 = r2 * r4; // s63 mul
|
|
r0 = r0 ^ r5; // s64 xor
|
|
r2 = r2 ^ r7; // s65 xor
|
|
r2 = r2 + r1 + ((((sel >> 6u) & 1u) != 0u) ? 0x8596687au : 0xd71c02ffu); // s66 add
|
|
r1 = rotl_imm(r1, 21u); // s67 rotl
|
|
r3 = r3 * r2; // s68 mul
|
|
r7 = r7 ^ __shfl_xor_sync(0xffffffffu, r5, 2); // s69 shfl
|
|
r3 = r3 * r2; // s70 mul
|
|
r0 = r2 * r5 + r0; // s71 mad
|
|
r6 = r6 + r7 + ((((sel >> 0u) & 1u) != 0u) ? 0x08ac5733u : 0x3c6fe15du); // s72 add
|
|
r3 = r3 ^ __shfl_xor_sync(0xffffffffu, r2, 8); // s73 shfl
|
|
r3 = r3 * r6; // s74 mul
|
|
r6 = r6 ^ r7; // s75 xor
|
|
r3 = r3 | r0; // s76 or
|
|
r2 = r2 + r4 + ((((sel >> 1u) & 1u) != 0u) ? 0xc100b495u : 0x295d5faeu); // s77 add
|
|
r3 = __umulhi(r3, r7); // s78 mulhi
|
|
r4 = r4 | r1; // s79 or
|
|
r4 = rotr_var(r4, r3); // s80 rotr
|
|
r4 = r4 + r3 + ((((sel >> 15u) & 1u) != 0u) ? 0x5cc59530u : 0x274a9221u); // s81 add
|
|
r7 = r7 + r3 + ((((sel >> 5u) & 1u) != 0u) ? 0x141479ecu : 0xc8651f8eu); // s82 add
|
|
r3 = rotr_var(r3, r6); // s83 rotr
|
|
r2 = r4 * r7 + r2; // s84 mad
|
|
r2 = r2 ^ __shfl_xor_sync(0xffffffffu, r5, 16); // s85 shfl
|
|
r3 = r3 ^ r2; // s86 xor
|
|
r5 = r5 ^ __shfl_xor_sync(0xffffffffu, r7, 2); // s87 shfl
|
|
r0 = r0 ^ r4; // s88 xor
|
|
r3 = r3 + r2 + ((((sel >> 18u) & 1u) != 0u) ? 0x883c0c92u : 0x53f915c2u); // s89 add
|
|
r7 = rotr_var(r7, r5); // s90 rotr
|
|
r3 = r3 + r1 + ((((sel >> 31u) & 1u) != 0u) ? 0x3cc5bd20u : 0xe2be00a5u); // s91 add
|
|
r7 = r7 ^ __shfl_xor_sync(0xffffffffu, r2, 1); // s92 shfl
|
|
r6 = rotr_var(r6, r2); // s93 rotr
|
|
r7 = rotl_imm(r7, 14u); // s94 rotl
|
|
r4 = r5 * r5 + r4; // s95 mad
|
|
r2 = r4 * r5 + r2; // s96 mad
|
|
r1 = r1 ^ r0; // s97 xor
|
|
r5 = r5 * r4; // s98 mul
|
|
r2 = r2 - r0; // s99 sub
|
|
r7 = rotl_imm(r7, 30u); // s100 rotl
|
|
r5 = r5 + r6 + ((((sel >> 17u) & 1u) != 0u) ? 0xbfd646cbu : 0x423fd9c9u); // s101 add
|
|
r7 = r7 + r6 + ((((sel >> 11u) & 1u) != 0u) ? 0x19b74a43u : 0x8d3c011du); // s102 add
|
|
r5 = rotl_imm(r5, 6u); // s103 rotl
|
|
r0 = r0 * r4; // s104 mul
|
|
r0 = r0 | r5; // s105 or
|
|
r0 = r0 + r1 + ((((sel >> 15u) & 1u) != 0u) ? 0x7dcb7e18u : 0x8ce14721u); // s106 add
|
|
r0 = rotl_imm(r0, 15u); // s107 rotl
|
|
r4 = r4 - r2; // s108 sub
|
|
r2 = __umulhi(r2, r0); // s109 mulhi
|
|
r1 = __umulhi(r1, r0); // s110 mulhi
|
|
r0 = r0 + r2 + ((((sel >> 28u) & 1u) != 0u) ? 0x3c179ad8u : 0x2d9fc6b9u); // s111 add
|
|
r2 = __umulhi(r2, r0); // s112 mulhi
|
|
r6 = r6 - r3; // s113 sub
|
|
r7 = r6 * r3 + r7; // s114 mad
|
|
r5 = rotr_var(r5, r1); // s115 rotr
|
|
r6 = rotr_var(r6, r4); // s116 rotr
|
|
r2 = r2 + r1 + ((((sel >> 22u) & 1u) != 0u) ? 0x47359729u : 0x502138e2u); // s117 add
|
|
r3 = r2 * r0 + r3; // s118 mad
|
|
r1 = r1 * r3; // s119 mul
|
|
r2 = r0 * r4 + r2; // s120 mad
|
|
r0 = r0 * r3; // s121 mul
|
|
r2 = __umulhi(r2, r0); // s122 mulhi
|
|
r5 = r5 * r6; // s123 mul
|
|
r4 = r2 * r4 + r4; // s124 mad
|
|
r4 = r4 ^ __shfl_xor_sync(0xffffffffu, r2, 2); // s125 shfl
|
|
r2 = r2 ^ r0; // s126 xor
|
|
r1 = rotl_imm(r1, 7u); // s127 rotl
|
|
r7 = r7 ^ __shfl_xor_sync(0xffffffffu, r2, 4); // s128 shfl
|
|
r6 = rotl_imm(r6, 17u); // s129 rotl
|
|
r2 = r2 + r5 + ((((sel >> 15u) & 1u) != 0u) ? 0x6c57e4e7u : 0x33dff776u); // s130 add
|
|
r0 = r0 | r3; // s131 or
|
|
r5 = rotr_var(r5, r0); // s132 rotr
|
|
r2 = r2 + r3 + ((((sel >> 30u) & 1u) != 0u) ? 0xe8212c0cu : 0x5db25b34u); // s133 add
|
|
r4 = r4 ^ r3; // s134 xor
|
|
r6 = r6 ^ r1; // s135 xor
|
|
r1 = r1 - r7; // s136 sub
|
|
r2 = r6 * r2 + r2; // s137 mad
|
|
r0 = __umulhi(r0, r5); // s138 mulhi
|
|
r2 = r2 + r7 + ((((sel >> 10u) & 1u) != 0u) ? 0xd44ff710u : 0x218d4090u); // s139 add
|
|
r1 = rotr_var(r1, r4); // s140 rotr
|
|
r3 = r3 ^ __shfl_xor_sync(0xffffffffu, r6, 1); // s141 shfl
|
|
r7 = r6 * r2 + r7; // s142 mad
|
|
r2 = r2 * r3; // s143 mul
|
|
r7 = r7 + r4 + ((((sel >> 29u) & 1u) != 0u) ? 0xb98942fau : 0xf4b1a8deu); // s144 add
|
|
r7 = rotl_imm(r7, 15u); // s145 rotl
|
|
r7 = r7 ^ r5; // s146 xor
|
|
r4 = r7 * r4 + r4; // s147 mad
|
|
r6 = rotr_var(r6, r5); // s148 rotr
|
|
r1 = r2 * r4 + r1; // s149 mad
|
|
r1 = r1 + r7 + ((((sel >> 17u) & 1u) != 0u) ? 0x0d48ba42u : 0x2bef10f2u); // s150 add
|
|
r5 = r5 ^ r4; // s151 xor
|
|
r7 = r7 + r0 + ((((sel >> 30u) & 1u) != 0u) ? 0xc98dea9cu : 0xdc5cc080u); // s152 add
|
|
r4 = r4 * r6; // s153 mul
|
|
r0 = r0 * r2; // s154 mul
|
|
r6 = r6 ^ r5; // s155 xor
|
|
r4 = rotr_var(r4, r2); // s156 rotr
|
|
r1 = rotl_imm(r1, 11u); // s157 rotl
|
|
r5 = r5 + r0 + ((((sel >> 11u) & 1u) != 0u) ? 0x6c752dcbu : 0x18197438u); // s158 add
|
|
r4 = r1 * r5 + r4; // s159 mad
|
|
r4 = rotl_imm(r4, 26u); // s160 rotl
|
|
r3 = r0 * r7 + r3; // s161 mad
|
|
r3 = rotr_var(r3, r1); // s162 rotr
|
|
r4 = r4 + r0 + ((((sel >> 3u) & 1u) != 0u) ? 0x8e481727u : 0xf1c46574u); // s163 add
|
|
r0 = __umulhi(r0, r4); // s164 mulhi
|
|
r2 = r2 + r6 + ((((sel >> 20u) & 1u) != 0u) ? 0x550ab406u : 0xac578137u); // s165 add
|
|
r0 = rotl_imm(r0, 13u); // s166 rotl
|
|
r3 = r3 + r1 + ((((sel >> 14u) & 1u) != 0u) ? 0xaebb5966u : 0xa33e6706u); // s167 add
|
|
r3 = r3 | r5; // s168 or
|
|
r6 = rotr_var(r6, r2); // s169 rotr
|
|
r4 = r4 ^ r6; // s170 xor
|
|
r6 = r6 - r1; // s171 sub
|
|
r7 = rotl_imm(r7, 22u); // s172 rotl
|
|
r5 = rotl_imm(r5, 15u); // s173 rotl
|
|
r7 = r7 ^ __shfl_xor_sync(0xffffffffu, r0, 8); // s174 shfl
|
|
r0 = r0 ^ r5; // s175 xor
|
|
r7 = rotl_imm(r7, 6u); // s176 rotl
|
|
r7 = r7 - r0; // s177 sub
|
|
r3 = rotl_imm(r3, 30u); // s178 rotl
|
|
r7 = r6 * r1 + r7; // s179 mad
|
|
r6 = rotl_imm(r6, 9u); // s180 rotl
|
|
r2 = r2 ^ r4; // s181 xor
|
|
r2 = r2 ^ r7; // s182 xor
|
|
r7 = r7 ^ r2; // s183 xor
|
|
r1 = r1 + r2 + ((((sel >> 21u) & 1u) != 0u) ? 0x5bb7550fu : 0xd94d55acu); // s184 add
|
|
r3 = r3 | r5; // s185 or
|
|
r6 = __umulhi(r6, r3); // s186 mulhi
|
|
r4 = r4 | r0; // s187 or
|
|
r7 = r7 ^ __shfl_xor_sync(0xffffffffu, r1, 2); // s188 shfl
|
|
r6 = r6 + r5 + ((((sel >> 6u) & 1u) != 0u) ? 0x2a354e2du : 0x89e747feu); // s189 add
|
|
r1 = r1 ^ r4; // s190 xor
|
|
r7 = __umulhi(r7, r4); // s191 mulhi
|
|
r2 = rotl_imm(r2, 26u); // s192 rotl
|
|
r5 = rotl_imm(r5, 8u); // s193 rotl
|
|
r4 = r4 ^ __shfl_xor_sync(0xffffffffu, r5, 16); // s194 shfl
|
|
r4 = r4 ^ r5; // s195 xor
|
|
r1 = r1 * r3; // s196 mul
|
|
r5 = r5 + r2 + ((((sel >> 7u) & 1u) != 0u) ? 0x10cfdc71u : 0xea03e8e7u); // s197 add
|
|
r7 = r7 + r5 + ((((sel >> 15u) & 1u) != 0u) ? 0x66e148d3u : 0xa7ee0102u); // s198 add
|
|
r5 = r5 - r2; // s199 sub
|
|
r5 = r5 ^ __shfl_xor_sync(0xffffffffu, r1, 2); // s200 shfl
|
|
r7 = r7 ^ __shfl_xor_sync(0xffffffffu, r1, 8); // s201 shfl
|
|
r4 = rotr_var(r4, r5); // s202 rotr
|
|
r7 = r7 ^ r5; // s203 xor
|
|
r7 = r7 ^ r0; // s204 xor
|
|
r6 = r6 + r2 + ((((sel >> 10u) & 1u) != 0u) ? 0x8558b619u : 0x8379a4deu); // s205 add
|
|
r4 = rotl_imm(r4, 13u); // s206 rotl
|
|
r1 = __umulhi(r1, r3); // s207 mulhi
|
|
r1 = r4 * r4 + r1; // s208 mad
|
|
r2 = r2 ^ __shfl_xor_sync(0xffffffffu, r0, 4); // s209 shfl
|
|
r7 = r7 + r0 + ((((sel >> 11u) & 1u) != 0u) ? 0xf4049c4cu : 0xaa8cb14eu); // s210 add
|
|
r7 = r7 ^ __shfl_xor_sync(0xffffffffu, r1, 16); // s211 shfl
|
|
r1 = r1 ^ r0; // s212 xor
|
|
r7 = rotl_imm(r7, 3u); // s213 rotl
|
|
r4 = rotr_var(r4, r7); // s214 rotr
|
|
r3 = r3 ^ __shfl_xor_sync(0xffffffffu, r2, 16); // s215 shfl
|
|
r5 = r5 | r1; // s216 or
|
|
r1 = r1 - r2; // s217 sub
|
|
r6 = r6 - r5; // s218 sub
|
|
r6 = rotl_imm(r6, 4u); // s219 rotl
|
|
r2 = __umulhi(r2, r0); // s220 mulhi
|
|
r2 = r2 | r0; // s221 or
|
|
r5 = r5 ^ __shfl_xor_sync(0xffffffffu, r2, 8); // s222 shfl
|
|
r5 = __umulhi(r5, r6); // s223 mulhi
|
|
r0 = r0 - r6; // s224 sub
|
|
r7 = rotl_imm(r7, 23u); // s225 rotl
|
|
r4 = r4 | r2; // s226 or
|
|
r2 = r2 * r4; // s227 mul
|
|
r3 = rotl_imm(r3, 12u); // s228 rotl
|
|
r0 = rotr_var(r0, r4); // s229 rotr
|
|
r0 = r0 + r6 + ((((sel >> 16u) & 1u) != 0u) ? 0x2a7fecb2u : 0x1d176220u); // s230 add
|
|
r1 = r1 ^ __shfl_xor_sync(0xffffffffu, r2, 4); // s231 shfl
|
|
r2 = r2 * r1; // s232 mul
|
|
r7 = r0 * r1 + r7; // s233 mad
|
|
r5 = rotl_imm(r5, 22u); // s234 rotl
|
|
r4 = rotr_var(r4, r6); // s235 rotr
|
|
r0 = r5 * r1 + r0; // s236 mad
|
|
r6 = r6 ^ r4; // s237 xor
|
|
r4 = r4 ^ r6; // s238 xor
|
|
r6 = rotl_imm(r6, 18u); // s239 rotl
|
|
r4 = r4 + r7 + ((((sel >> 25u) & 1u) != 0u) ? 0xadce39f3u : 0x17dafb4du); // s240 add
|
|
r0 = r0 - r7; // s241 sub
|
|
r1 = rotr_var(r1, r6); // s242 rotr
|
|
r3 = r3 + r0 + ((((sel >> 29u) & 1u) != 0u) ? 0x0e7033b6u : 0xf2f77d26u); // s243 add
|
|
r2 = r7 * r4 + r2; // s244 mad
|
|
r4 = r0 * r6 + r4; // s245 mad
|
|
r5 = r5 * r7; // s246 mul
|
|
r3 = r3 + r5 + ((((sel >> 31u) & 1u) != 0u) ? 0x7b2ff6b7u : 0xfed2da4eu); // s247 add
|
|
r0 = r0 + r7 + ((((sel >> 0u) & 1u) != 0u) ? 0xb5fad7b1u : 0x03b2891cu); // s248 add
|
|
r5 = __umulhi(r5, r4); // s249 mulhi
|
|
r5 = r5 + r2 + ((((sel >> 11u) & 1u) != 0u) ? 0xa7e71c2fu : 0x042cd6e3u); // s250 add
|
|
r6 = r6 ^ __shfl_xor_sync(0xffffffffu, r1, 1); // s251 shfl
|
|
r6 = r6 ^ r1; // s252 xor
|
|
r2 = r2 * r5; // s253 mul
|
|
r0 = r0 + r6 + ((((sel >> 16u) & 1u) != 0u) ? 0xb83d78deu : 0x784a302bu); // s254 add
|
|
r2 = r2 - r4; // s255 sub
|
|
}
|
|
}
|
|
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);
|
|
}
|