541 lines
32 KiB
Common Lisp
541 lines
32 KiB
Common Lisp
// Generated by igneum-pow export (generator v2) for seed "igneum-epoch/edc4fa844da9dc98d37e965176f6558a31560e40502ab3ae5491b21aaaabfb07/day/69676e65756d2d6461792ffa50000000000000". Do not edit by hand.
|
|
// OpenCL C twin of the Metal kernel for the same seed (see proto-opencl/README.md, WAVEFRONT.md and program.metal).
|
|
// Built from source at runtime by proto-opencl/host.c, which passes these defines:
|
|
// IGNEUM_GROUP work-group size of igneum_hash, a multiple of 32 (default 32: one work-group = one 32-lane unit)
|
|
// IGNEUM_EXCHANGE 0 = local-memory exchange with a barrier (any device, any wave width; the default)
|
|
// 1 = sub_group_shuffle_xor (cl_khr_subgroup_shuffle), only with IGNEUM_GROUP 32 and a sub-group size of exactly 32
|
|
// 2 = intel_sub_group_shuffle_xor (cl_intel_subgroups), same condition
|
|
// The verification unit is always 32 lanes. A 64-wide hardware wave (AMD GCN/CDNA, RDNA in wave64) runs two units;
|
|
// the exchange masks are 1, 2, 4, 8, 16, so every partner lane lies inside the lane's own aligned run of 32.
|
|
#ifndef IGNEUM_GROUP
|
|
#define IGNEUM_GROUP 32
|
|
#endif
|
|
#ifndef IGNEUM_EXCHANGE
|
|
#define IGNEUM_EXCHANGE 0
|
|
#endif
|
|
#ifdef __OPENCL_VERSION__
|
|
#define IGNEUM_KERNEL_HASH __kernel __attribute__((reqd_work_group_size(IGNEUM_GROUP, 1, 1)))
|
|
#define IGNEUM_LOCAL_WORDS(name, n) __local uint name[n]
|
|
#if IGNEUM_EXCHANGE == 1
|
|
#ifdef cl_khr_subgroups
|
|
#pragma OPENCL EXTENSION cl_khr_subgroups : enable
|
|
#endif
|
|
#ifdef cl_khr_subgroup_shuffle
|
|
#pragma OPENCL EXTENSION cl_khr_subgroup_shuffle : enable
|
|
#endif
|
|
#elif IGNEUM_EXCHANGE == 2
|
|
#pragma OPENCL EXTENSION cl_intel_subgroups : enable
|
|
#endif
|
|
#else
|
|
// Not an OpenCL compiler: proto-opencl/emu compiles this file as C++ and supplies the built-ins and these two macros.
|
|
#include "emu_opencl.h"
|
|
#endif
|
|
|
|
#if IGNEUM_EXCHANGE == 1
|
|
#define IGNEUM_SHFL_XOR(dst, a, m) dst = sub_group_shuffle_xor((a), (uint)(m))
|
|
#define IGNEUM_BCAST0(dst, a) dst = sub_group_broadcast((a), 0u)
|
|
#elif IGNEUM_EXCHANGE == 2
|
|
#define IGNEUM_SHFL_XOR(dst, a, m) dst = intel_sub_group_shuffle_xor((a), (uint)(m))
|
|
#define IGNEUM_BCAST0(dst, a) dst = sub_group_broadcast((a), 0u)
|
|
#else
|
|
// Local-memory exchange. Two buffers of IGNEUM_GROUP words alternate (xk counts exchanges), so one barrier per
|
|
// exchange is enough: a lane can only overwrite buffer b at exchange k+2 after passing barrier k+1, and every lane
|
|
// reaches barrier k+1 only after its read of buffer b at exchange k. The partner lid ^ m stays inside the lane's
|
|
// aligned run of 32 because m < 32. Control flow is uniform, so every work-item reaches every barrier.
|
|
#define IGNEUM_SHFL_XOR(dst, a, m) { xch[(xk & 1u) * IGNEUM_GROUP + lid] = (a); barrier(CLK_LOCAL_MEM_FENCE); dst = xch[(xk & 1u) * IGNEUM_GROUP + (lid ^ (uint)(m))]; xk += 1u; }
|
|
#define IGNEUM_BCAST0(dst, a) { xch[(xk & 1u) * IGNEUM_GROUP + lid] = (a); barrier(CLK_LOCAL_MEM_FENCE); dst = xch[(xk & 1u) * IGNEUM_GROUP + (lid & ~31u)]; xk += 1u; }
|
|
#endif
|
|
|
|
static inline uint splitmix32(uint 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. OpenCL rotate() rotates left by n modulo 32.
|
|
static inline uint rotl_imm(uint x, uint n) { return rotate(x, n); }
|
|
// Right rotation by n modulo 32 as a left rotation by (32 - n) modulo 32; n == 0 gives x.
|
|
static inline uint rotr_var(uint x, uint n) { return rotate(x, (0u - n) & 31u); }
|
|
static inline uint ds_elem(uint i, uint d0, uint d1) {
|
|
uint 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 core (MEMHARD.md). Cache: 2^26 words in 2^16 segments of 64 chained ChaCha12 lines.
|
|
// Item: 8 rounds of 8 x seed-parameterised mixer + one 64-byte cache read, then 8 x final mixer (class v3, mixer multiplier 8,
|
|
// docs/plans/mixer-x4.md: the round key of application j of round r is 0x9E3779B9 * (r * m + j + 1)). All parameters are literals.
|
|
#define MH_CACHE_LINE_MASK 0x003fffffu
|
|
#define MH_SEGMENT_LINES 64u
|
|
#define MH_QR(a, b, c, d, r1, r2, r3, r4) { a += b; d ^= a; d = mh_rotl(d, r1); c += d; b ^= c; b = mh_rotl(b, r2); a += b; d ^= a; d = mh_rotl(d, r3); c += d; b ^= c; b = mh_rotl(b, r4); }
|
|
static inline uint mh_rotl(uint x, uint n) { return (x << n) | (x >> (32u - n)); } // n in 1..31 at every call site
|
|
|
|
// y = ChaCha12 core(x) + x
|
|
static inline void mh_chacha_block(const uint* x, uint* y) {
|
|
for (uint i = 0u; i < 16u; ++i) y[i] = x[i];
|
|
for (uint r = 0u; r < 6u; ++r) {
|
|
MH_QR(y[0], y[4], y[8], y[12], 16u, 12u, 8u, 7u) MH_QR(y[1], y[5], y[9], y[13], 16u, 12u, 8u, 7u)
|
|
MH_QR(y[2], y[6], y[10], y[14], 16u, 12u, 8u, 7u) MH_QR(y[3], y[7], y[11], y[15], 16u, 12u, 8u, 7u)
|
|
MH_QR(y[0], y[5], y[10], y[15], 16u, 12u, 8u, 7u) MH_QR(y[1], y[6], y[11], y[12], 16u, 12u, 8u, 7u)
|
|
MH_QR(y[2], y[7], y[8], y[13], 16u, 12u, 8u, 7u) MH_QR(y[3], y[4], y[9], y[14], 16u, 12u, 8u, 7u)
|
|
}
|
|
for (uint i = 0u; i < 16u; ++i) y[i] += x[i];
|
|
}
|
|
|
|
// One cache segment: 64 chained lines written at cache[seg * 1024]. in_j = prev ^ (sigma || K || seg || j || tag), prev_0 = 0.
|
|
static inline void mh_cache_segment(__global uint* cache, uint seg) {
|
|
uint prev[16]; uint x[16]; uint y[16];
|
|
for (uint i = 0u; i < 16u; ++i) prev[i] = 0u;
|
|
for (uint j = 0u; j < MH_SEGMENT_LINES; ++j) {
|
|
x[0] = 0x61707865u ^ prev[0]; x[1] = 0x3320646eu ^ prev[1]; x[2] = 0x79622d32u ^ prev[2]; x[3] = 0x6b206574u ^ prev[3];
|
|
x[4] = 0xceed56d7u ^ prev[4];
|
|
x[5] = 0x9ba270d2u ^ prev[5];
|
|
x[6] = 0x82caab2du ^ prev[6];
|
|
x[7] = 0x81ebce0eu ^ prev[7];
|
|
x[8] = 0x12b6ecf1u ^ prev[8];
|
|
x[9] = 0xd0f3fd7cu ^ prev[9];
|
|
x[10] = 0xd872eefeu ^ prev[10];
|
|
x[11] = 0xc158c7bdu ^ prev[11];
|
|
x[12] = seg ^ prev[12]; x[13] = j ^ prev[13]; x[14] = 0x49676e65u ^ prev[14]; x[15] = 0x756d4d48u ^ prev[15];
|
|
mh_chacha_block(x, y);
|
|
__global uint* line = cache + ((seg * MH_SEGMENT_LINES + j) * 16u);
|
|
for (uint i = 0u; i < 16u; ++i) { line[i] = y[i]; prev[i] = y[i]; }
|
|
}
|
|
}
|
|
|
|
// M_r: per word (s ^ (RC + rk)) * MUL, then a column round and a diagonal round with the seed-drawn rotations.
|
|
static inline void mh_mixer(uint* s, uint rk) {
|
|
s[0] = (s[0] ^ (0xc6892460u + rk)) * 0xf351d601u;
|
|
s[1] = (s[1] ^ (0x25b7228au + rk)) * 0xa3bb398fu;
|
|
s[2] = (s[2] ^ (0xcd515004u + rk)) * 0xb5a09e35u;
|
|
s[3] = (s[3] ^ (0x2846527au + rk)) * 0x7509c9c1u;
|
|
s[4] = (s[4] ^ (0xa6324241u + rk)) * 0x6bbf31e9u;
|
|
s[5] = (s[5] ^ (0x36e3ec53u + rk)) * 0xfc849a79u;
|
|
s[6] = (s[6] ^ (0x82961bacu + rk)) * 0xded91851u;
|
|
s[7] = (s[7] ^ (0x0f97ba7du + rk)) * 0x8d9113d1u;
|
|
s[8] = (s[8] ^ (0xb6f921a9u + rk)) * 0x0ff15225u;
|
|
s[9] = (s[9] ^ (0x3ada24e5u + rk)) * 0x3a5bdd41u;
|
|
s[10] = (s[10] ^ (0xde20ab91u + rk)) * 0xab533435u;
|
|
s[11] = (s[11] ^ (0x5378eeb2u + rk)) * 0xe1c55ad5u;
|
|
s[12] = (s[12] ^ (0x7d161662u + rk)) * 0xe6d3bd0du;
|
|
s[13] = (s[13] ^ (0x89353cc1u + rk)) * 0x9d9ffbbdu;
|
|
s[14] = (s[14] ^ (0xb1aa03a2u + rk)) * 0xbb2a3cf3u;
|
|
s[15] = (s[15] ^ (0x788acae6u + rk)) * 0x50a7c08du;
|
|
MH_QR(s[0], s[4], s[8], s[12], 17u, 12u, 20u, 23u) MH_QR(s[1], s[5], s[9], s[13], 17u, 12u, 20u, 23u)
|
|
MH_QR(s[2], s[6], s[10], s[14], 17u, 12u, 20u, 23u) MH_QR(s[3], s[7], s[11], s[15], 17u, 12u, 20u, 23u)
|
|
MH_QR(s[0], s[5], s[10], s[15], 7u, 3u, 27u, 16u) MH_QR(s[1], s[6], s[11], s[12], 7u, 3u, 27u, 16u)
|
|
MH_QR(s[2], s[7], s[8], s[13], 7u, 3u, 27u, 16u) MH_QR(s[3], s[4], s[9], s[14], 7u, 3u, 27u, 16u)
|
|
}
|
|
|
|
// Item t: 16 words. s = (K, t * MUL[i] + RC[i]); 8 rounds of 8 x mixer + cache line s[0] & mask; 8 x final mixer.
|
|
static inline void mh_item(__global const uint* cache, uint t, uint* s) {
|
|
s[0] = 0xceed56d7u;
|
|
s[1] = 0x9ba270d2u;
|
|
s[2] = 0x82caab2du;
|
|
s[3] = 0x81ebce0eu;
|
|
s[4] = 0x12b6ecf1u;
|
|
s[5] = 0xd0f3fd7cu;
|
|
s[6] = 0xd872eefeu;
|
|
s[7] = 0xc158c7bdu;
|
|
s[8] = t * 0xf351d601u + 0xc6892460u;
|
|
s[9] = t * 0xa3bb398fu + 0x25b7228au;
|
|
s[10] = t * 0xb5a09e35u + 0xcd515004u;
|
|
s[11] = t * 0x7509c9c1u + 0x2846527au;
|
|
s[12] = t * 0x6bbf31e9u + 0xa6324241u;
|
|
s[13] = t * 0xfc849a79u + 0x36e3ec53u;
|
|
s[14] = t * 0xded91851u + 0x82961bacu;
|
|
s[15] = t * 0x8d9113d1u + 0x0f97ba7du;
|
|
for (uint r = 0u; r < 8u; ++r) {
|
|
for (uint j = 0u; j < 8u; ++j) mh_mixer(s, 0x9E3779B9u * (r * 8u + j + 1u));
|
|
__global const uint* line = cache + ((s[0] & MH_CACHE_LINE_MASK) * 16u);
|
|
for (uint i = 0u; i < 16u; ++i) s[i] ^= line[i];
|
|
}
|
|
for (uint j = 0u; j < 8u; ++j) mh_mixer(s, 0x9E3779B9u * (64u + j + 1u));
|
|
}
|
|
// Era layout (docs/plans/era-layout.md 1.2): dataset word w holds word mh_j(w) of item mh_t(w); j's bits sit at positions 0 2 12 13 of w.
|
|
static inline uint mh_j(uint w) { return ((w >> 0u) & 1u) | (((w >> 2u) & 1u) << 1) | (((w >> 12u) & 1u) << 2) | (((w >> 13u) & 1u) << 3); }
|
|
static inline uint mh_t(uint w) { w = (w & 0x00001fffu) | ((w >> 14u) << 13u); w = (w & 0x00000fffu) | ((w >> 13u) << 12u); w = (w & 0x00000003u) | ((w >> 3u) << 2u); w = (w & 0x00000000u) | ((w >> 1u) << 0u); return w; }
|
|
static inline uint mh_addr(uint t, uint j) { uint w = t; w = ((w >> 0u) << 1u) | (w & 0x00000000u) | (((j >> 0u) & 1u) << 0u); w = ((w >> 2u) << 3u) | (w & 0x00000003u) | (((j >> 1u) & 1u) << 2u); w = ((w >> 12u) << 13u) | (w & 0x00000fffu) | (((j >> 2u) & 1u) << 12u); w = ((w >> 13u) << 14u) | (w & 0x00001fffu) | (((j >> 3u) & 1u) << 13u); return w; }
|
|
// dataset[w] without the dataset: derive item mh_t(w) and take word mh_j(w).
|
|
static inline uint mh_word(__global const uint* cache, uint w) { uint s[16]; mh_item(cache, mh_t(w), s); return s[mh_j(w)]; }
|
|
|
|
// Memory-hard dataset (MEMHARD.md). One work-item per cache segment; one work-item per 64-byte dataset item.
|
|
// The same constants as memhard.h in this pack (one emitter, three dialects).
|
|
__kernel void igneum_cache_fill(__global uint* cache, uint nSegments) {
|
|
uint seg = (uint)get_global_id(0);
|
|
if (seg < nSegments) mh_cache_segment(cache, seg);
|
|
}
|
|
__kernel void igneum_build(__global uint* ds, __global const uint* cache, uint nItems) {
|
|
uint t = (uint)get_global_id(0);
|
|
if (t < nItems) {
|
|
uint s[16];
|
|
mh_item(cache, t, s);
|
|
for (uint i = 0u; i < 16u; ++i) ds[(ulong)mh_addr(t, i)] = s[i];
|
|
}
|
|
}
|
|
|
|
// One hash per work-item. IGNEUM_GROUP is a multiple of 32; lane = lid & 31 and every exchange stays inside the
|
|
// lane's own aligned run of 32 work-items, exactly like simd_shuffle_xor inside a 32-wide Metal SIMD group and
|
|
// __shfl_xor_sync inside a CUDA warp. Control flow is uniform (no branches at all).
|
|
IGNEUM_KERNEL_HASH void igneum_hash(__global const uint* ds, __global ulong* out, uint baseNonce, uint mask) {
|
|
uint gid = (uint)get_global_id(0);
|
|
uint lid = (uint)get_local_id(0);
|
|
uint nonce = baseNonce + gid;
|
|
uint r0, r1, r2, r3, r4, r5, r6, r7;
|
|
#if IGNEUM_EXCHANGE == 0
|
|
IGNEUM_LOCAL_WORDS(xch, 2 * IGNEUM_GROUP);
|
|
uint xk = 0u;
|
|
#else
|
|
(void)lid;
|
|
#endif
|
|
{ uint x = nonce ^ 0x667d0fbdu; x += 0x9e3779b9u; x = splitmix32(x); r0 = x ^ 0x7b8e5963u; } // SEEDW[0], 0x9e3779b9u * 1u, SEEDW[1]
|
|
{ uint x = nonce ^ 0x7b8e5963u; x += 0x3c6ef372u; x = splitmix32(x); r1 = x ^ 0x31c67e5eu; } // SEEDW[1], 0x9e3779b9u * 2u, SEEDW[2]
|
|
{ uint x = nonce ^ 0x31c67e5eu; x += 0xdaa66d2bu; x = splitmix32(x); r2 = x ^ 0x4529ddc6u; } // SEEDW[2], 0x9e3779b9u * 3u, SEEDW[3]
|
|
{ uint x = nonce ^ 0x4529ddc6u; x += 0x78dde6e4u; x = splitmix32(x); r3 = x ^ 0xef19d6d8u; } // SEEDW[3], 0x9e3779b9u * 4u, SEEDW[4]
|
|
{ uint x = nonce ^ 0xef19d6d8u; x += 0x1715609du; x = splitmix32(x); r4 = x ^ 0xaccf6211u; } // SEEDW[4], 0x9e3779b9u * 5u, SEEDW[5]
|
|
{ uint x = nonce ^ 0xaccf6211u; x += 0xb54cda56u; x = splitmix32(x); r5 = x ^ 0xda0aed32u; } // SEEDW[5], 0x9e3779b9u * 6u, SEEDW[6]
|
|
{ uint x = nonce ^ 0xda0aed32u; x += 0x5384540fu; x = splitmix32(x); r6 = x ^ 0xabc6df31u; } // SEEDW[6], 0x9e3779b9u * 7u, SEEDW[7]
|
|
{ uint x = nonce ^ 0xabc6df31u; x += 0xf1bbcdc8u; x = splitmix32(x); r7 = x ^ 0x667d0fbdu; } // SEEDW[7], 0x9e3779b9u * 8u, SEEDW[0]
|
|
|
|
for (uint it = 0u; it < 8u; ++it) {
|
|
uint sel = r0;
|
|
r4 = r4 + r5 + ((((sel >> 13u) & 1u) != 0u) ? 0x5810667au : 0xea86e152u); // 0 add
|
|
r7 = r7 ^ r0; // 1 xor
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r6, 1u); r3 = r3 ^ t_; } // 2 shfl
|
|
r4 = r4 - r1; // 3 sub
|
|
r2 = r0 * r4 + r2; // 4 mad
|
|
r4 = rotl_imm(r4, 9u); // 5 rotl
|
|
r0 = rotr_var(r0, r2); // 6 rotr
|
|
r6 = r6 ^ ds[((rotl_imm(r7 * 0x9ad30d99u, 29u) & 0x03ffffffu) | 0x04000000u) & mask]; // 7 load
|
|
r1 = r1 ^ ds[((rotl_imm(r4 * 0x9ad30d99u, 29u) & 0x07ffffffu) | 0x08000000u) & mask]; // 8 load
|
|
r1 = r1 ^ ds[((rotl_imm(r2 * 0x9ad30d99u, 29u) & 0x07ffffffu) | 0x08000000u) & mask]; // 9 load
|
|
r7 = r7 ^ ds[((rotl_imm(r0 * 0x9ad30d99u, 29u) & 0x07ffffffu) | 0x08000000u) & mask]; // 10 load
|
|
r7 = r7 ^ ds[((rotl_imm(r1 * 0x9ad30d99u, 29u) & 0x0fffffffu) | 0x00000000u) & mask]; // 11 load
|
|
r5 = r5 * r4; // 12 mul
|
|
r7 = r7 ^ ds[((rotl_imm(r3 * 0x9ad30d99u, 29u) & 0x07ffffffu) | 0x08000000u) & mask]; // 13 load
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r5, 16u); r6 = r6 ^ t_; } // 14 shfl
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r5, 1u); r3 = r3 ^ t_; } // 15 shfl
|
|
r2 = r2 | r7; // 16 or
|
|
r0 = rotr_var(r0, r6); // 17 rotr
|
|
r7 = r7 + r5 + ((((sel >> 12u) & 1u) != 0u) ? 0xb9e3577eu : 0xf66e7017u); // 18 add
|
|
r7 = rotl_imm(r7, 24u); // 19 rotl
|
|
r6 = r6 + r7 + ((((sel >> 23u) & 1u) != 0u) ? 0x8c9f0ef8u : 0x52334d12u); // 20 add
|
|
r2 = r2 | r6; // 21 or
|
|
r6 = r5 * r4 + r6; // 22 mad
|
|
r1 = r1 | r0; // 23 or
|
|
r6 = r6 ^ r1; // 24 xor
|
|
r2 = r2 + r6 + ((((sel >> 17u) & 1u) != 0u) ? 0x9cec0e12u : 0x659fc3d3u); // 25 add
|
|
r7 = rotr_var(r7, r0); // 26 rotr
|
|
r4 = r5 * r7 + r4; // 27 mad
|
|
r3 = r2 * r4 + r3; // 28 mad
|
|
r1 = r1 ^ ds[((rotl_imm(r6 * 0x9ad30d99u, 29u) & 0x0fffffffu) | 0x00000000u) & mask]; // 29 load
|
|
r2 = r2 ^ ds[((rotl_imm(r3 * 0x9ad30d99u, 29u) & 0x03ffffffu) | 0x08000000u) & mask]; // 30 load
|
|
r1 = r1 ^ ds[((rotl_imm(r4 * 0x9ad30d99u, 29u) & 0x07ffffffu) | 0x08000000u) & mask]; // 31 load
|
|
r7 = r7 + r2 + ((((sel >> 16u) & 1u) != 0u) ? 0xe403240eu : 0x070888a8u); // 32 add
|
|
r2 = r2 + r0 + ((((sel >> 28u) & 1u) != 0u) ? 0x29701828u : 0xf2e46d55u); // 33 add
|
|
r2 = r2 + r3 + ((((sel >> 14u) & 1u) != 0u) ? 0x343b7aeeu : 0x58f75b87u); // 34 add
|
|
r2 = mul_hi(r2, r5); // 35 mulhi
|
|
r4 = r4 ^ r2; // 36 xor
|
|
r6 = r6 * r5; // 37 mul
|
|
r7 = r7 ^ r0; // 38 xor
|
|
r7 = r7 + r2 + ((((sel >> 19u) & 1u) != 0u) ? 0x32bbd117u : 0xb8180e9du); // 39 add
|
|
r2 = rotr_var(r2, r3); // 40 rotr
|
|
r7 = r7 - r0; // 41 sub
|
|
r4 = r4 + r3 + ((((sel >> 4u) & 1u) != 0u) ? 0x6df7aed4u : 0x6ced15b7u); // 42 add
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r3, 4u); r7 = r7 ^ t_; } // 43 shfl
|
|
r0 = r0 ^ ds[((rotl_imm(r2 * 0x9ad30d99u, 29u) & 0x07ffffffu) | 0x08000000u) & mask]; // 44 load
|
|
r1 = r1 + r6 + ((((sel >> 14u) & 1u) != 0u) ? 0x83e825bfu : 0xe09f54e9u); // 45 add
|
|
r3 = r3 ^ ds[((rotl_imm(r0 * 0x9ad30d99u, 29u) & 0x03ffffffu) | 0x00000000u) & mask]; // 46 load
|
|
r6 = r6 ^ ds[((rotl_imm(r4 * 0x9ad30d99u, 29u) & 0x0fffffffu) | 0x00000000u) & mask]; // 47 load
|
|
r4 = mul_hi(r4, r2); // 48 mulhi
|
|
r5 = r5 + r0 + ((((sel >> 7u) & 1u) != 0u) ? 0xf572bdb9u : 0xa8bae6dfu); // 49 add
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r7, 4u); r0 = r0 ^ t_; } // 50 shfl
|
|
r6 = r6 + r0 + ((((sel >> 19u) & 1u) != 0u) ? 0x11e17c61u : 0x383b9260u); // 51 add
|
|
r5 = r5 ^ ds[((rotl_imm(r0 * 0x9ad30d99u, 29u) & 0x0fffffffu) | 0x00000000u) & mask]; // 52 load
|
|
r6 = rotl_imm(r6, 12u); // 53 rotl
|
|
r3 = rotl_imm(r3, 12u); // 54 rotl
|
|
r2 = r2 + r1 + ((((sel >> 10u) & 1u) != 0u) ? 0x18d67dbbu : 0xac6be8e3u); // 55 add
|
|
r1 = r1 ^ ds[((rotl_imm(r6 * 0x9ad30d99u, 29u) & 0x0fffffffu) | 0x00000000u) & mask]; // 56 load
|
|
r5 = r5 + r0 + ((((sel >> 12u) & 1u) != 0u) ? 0xa75cd60du : 0xf03673feu); // 57 add
|
|
r5 = r5 ^ ds[((rotl_imm(r7 * 0x9ad30d99u, 29u) & 0x03ffffffu) | 0x00000000u) & mask]; // 58 load
|
|
r1 = r1 + r4 + ((((sel >> 7u) & 1u) != 0u) ? 0x8dfb96bbu : 0xdecd4794u); // 59 add
|
|
r3 = mul_hi(r3, r2); // 60 mulhi
|
|
r6 = r6 | r4; // 61 or
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r4, 8u); r5 = r5 ^ t_; } // 62 shfl
|
|
r3 = r3 ^ ds[((rotl_imm(r2 * 0x9ad30d99u, 29u) & 0x07ffffffu) | 0x08000000u) & mask]; // 63 load
|
|
// latency-shadow block (Counter ASIC 3.0 item 8): 256 ALU instructions x 27 passes after instruction 63, no load
|
|
for (uint sh = 0u; sh < 27u; ++sh) {
|
|
r7 = r7 * r3; // s0 mul
|
|
r7 = r7 + r4 + ((((sel >> 16u) & 1u) != 0u) ? 0x06575fd2u : 0xecb44e9cu); // s1 add
|
|
r5 = mul_hi(r5, r0); // s2 mulhi
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r5, 2u); r4 = r4 ^ t_; } // s3 shfl
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r1, 1u); r4 = r4 ^ t_; } // s4 shfl
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r5, 8u); r4 = r4 ^ t_; } // s5 shfl
|
|
r6 = r6 - r1; // s6 sub
|
|
r3 = r3 | r4; // s7 or
|
|
r0 = r4 * r7 + r0; // s8 mad
|
|
r3 = r3 - r4; // s9 sub
|
|
r6 = r6 * r3; // s10 mul
|
|
r5 = mul_hi(r5, r0); // s11 mulhi
|
|
r0 = r4 * r5 + r0; // s12 mad
|
|
r3 = r3 + r7 + ((((sel >> 29u) & 1u) != 0u) ? 0x41da8352u : 0x78295146u); // s13 add
|
|
r7 = r7 ^ r2; // s14 xor
|
|
r1 = r1 + r4 + ((((sel >> 12u) & 1u) != 0u) ? 0x491bea93u : 0x1126a483u); // s15 add
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r2, 8u); r1 = r1 ^ t_; } // s16 shfl
|
|
r5 = rotr_var(r5, r0); // s17 rotr
|
|
r2 = r2 - r1; // s18 sub
|
|
r3 = mul_hi(r3, r2); // s19 mulhi
|
|
r0 = r0 ^ r4; // s20 xor
|
|
r2 = r2 + r3 + ((((sel >> 31u) & 1u) != 0u) ? 0xd0df3d0au : 0x7289ac7cu); // s21 add
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r7, 2u); r2 = r2 ^ t_; } // s22 shfl
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r0, 8u); r1 = r1 ^ t_; } // s23 shfl
|
|
r1 = r1 - r2; // s24 sub
|
|
r7 = r3 * r1 + r7; // s25 mad
|
|
r2 = r2 ^ r6; // s26 xor
|
|
r4 = mul_hi(r4, r2); // s27 mulhi
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r2, 2u); r1 = r1 ^ t_; } // s28 shfl
|
|
r2 = r2 - r5; // s29 sub
|
|
r7 = mul_hi(r7, r6); // s30 mulhi
|
|
r2 = rotr_var(r2, r7); // s31 rotr
|
|
r6 = rotl_imm(r6, 26u); // s32 rotl
|
|
r0 = r0 * r7; // s33 mul
|
|
r7 = r7 * r4; // s34 mul
|
|
r6 = r0 * r1 + r6; // s35 mad
|
|
r4 = r4 + r2 + ((((sel >> 4u) & 1u) != 0u) ? 0xb3e87396u : 0xfe2b1c7du); // s36 add
|
|
r7 = rotr_var(r7, r4); // s37 rotr
|
|
r5 = r2 * r7 + r5; // s38 mad
|
|
r6 = r6 + r2 + ((((sel >> 16u) & 1u) != 0u) ? 0x31a906adu : 0xe036a560u); // s39 add
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r6, 4u); r3 = r3 ^ t_; } // s40 shfl
|
|
r5 = r5 ^ r0; // s41 xor
|
|
r5 = r5 ^ r3; // s42 xor
|
|
r6 = rotl_imm(r6, 31u); // s43 rotl
|
|
r0 = rotl_imm(r0, 6u); // s44 rotl
|
|
r2 = r0 * r6 + r2; // s45 mad
|
|
r0 = r6 * r0 + r0; // s46 mad
|
|
r5 = r5 ^ r2; // s47 xor
|
|
r1 = r1 + r5 + ((((sel >> 27u) & 1u) != 0u) ? 0xc839ad2eu : 0xd802a3efu); // s48 add
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r1, 2u); r0 = r0 ^ t_; } // s49 shfl
|
|
r6 = r6 + r0 + ((((sel >> 18u) & 1u) != 0u) ? 0xe9d06965u : 0x8e71c4c0u); // s50 add
|
|
r1 = rotl_imm(r1, 3u); // s51 rotl
|
|
r6 = r6 ^ r2; // s52 xor
|
|
r5 = r5 ^ r6; // s53 xor
|
|
r2 = r2 * r6; // s54 mul
|
|
r0 = mul_hi(r0, r2); // s55 mulhi
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r3, 1u); r6 = r6 ^ t_; } // s56 shfl
|
|
r1 = r1 + r2 + ((((sel >> 21u) & 1u) != 0u) ? 0x5b84a832u : 0xb0eb7d4eu); // s57 add
|
|
r0 = rotr_var(r0, r1); // s58 rotr
|
|
r6 = r6 + r4 + ((((sel >> 8u) & 1u) != 0u) ? 0x4e1a16c6u : 0xd35e37c1u); // s59 add
|
|
r0 = r0 | r2; // s60 or
|
|
r5 = rotr_var(r5, r3); // s61 rotr
|
|
r4 = r4 - r2; // s62 sub
|
|
r0 = r0 + r1 + ((((sel >> 27u) & 1u) != 0u) ? 0xec29413au : 0x16221227u); // s63 add
|
|
r6 = rotl_imm(r6, 20u); // s64 rotl
|
|
r1 = r1 ^ r2; // s65 xor
|
|
r6 = rotl_imm(r6, 23u); // s66 rotl
|
|
r1 = r1 | r4; // s67 or
|
|
r7 = r4 * r0 + r7; // s68 mad
|
|
r1 = r1 | r5; // s69 or
|
|
r7 = r7 ^ r4; // s70 xor
|
|
r2 = r2 * r3; // s71 mul
|
|
r0 = r0 * r2; // s72 mul
|
|
r7 = r7 + r6 + ((((sel >> 4u) & 1u) != 0u) ? 0x85c0f694u : 0x5708524au); // s73 add
|
|
r3 = r7 * r0 + r3; // s74 mad
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r2, 8u); r0 = r0 ^ t_; } // s75 shfl
|
|
r5 = r5 - r7; // s76 sub
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r1, 2u); r5 = r5 ^ t_; } // s77 shfl
|
|
r5 = r5 + r0 + ((((sel >> 26u) & 1u) != 0u) ? 0xad4f292bu : 0xc4195db9u); // s78 add
|
|
r5 = r5 * r6; // s79 mul
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r7, 2u); r6 = r6 ^ t_; } // s80 shfl
|
|
r5 = r5 * r6; // s81 mul
|
|
r7 = r7 | r3; // s82 or
|
|
r0 = r4 * r2 + r0; // s83 mad
|
|
r4 = rotl_imm(r4, 12u); // s84 rotl
|
|
r3 = r3 * r4; // s85 mul
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r2, 4u); r0 = r0 ^ t_; } // s86 shfl
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r7, 4u); r2 = r2 ^ t_; } // s87 shfl
|
|
r1 = r1 ^ r3; // s88 xor
|
|
r5 = r5 ^ r1; // s89 xor
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r1, 16u); r6 = r6 ^ t_; } // s90 shfl
|
|
r5 = r5 + r0 + ((((sel >> 7u) & 1u) != 0u) ? 0xb41704ddu : 0x5570a07eu); // s91 add
|
|
r0 = mul_hi(r0, r1); // s92 mulhi
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r0, 8u); r6 = r6 ^ t_; } // s93 shfl
|
|
r3 = r3 + r6 + ((((sel >> 18u) & 1u) != 0u) ? 0x4674baa3u : 0xcd1be982u); // s94 add
|
|
r4 = rotl_imm(r4, 24u); // s95 rotl
|
|
r3 = r7 * r4 + r3; // s96 mad
|
|
r6 = r6 - r3; // s97 sub
|
|
r1 = r5 * r6 + r1; // s98 mad
|
|
r6 = rotr_var(r6, r2); // s99 rotr
|
|
r5 = r5 ^ r1; // s100 xor
|
|
r3 = r3 + r7 + ((((sel >> 7u) & 1u) != 0u) ? 0x3d25f873u : 0x60f87347u); // s101 add
|
|
r3 = r3 - r5; // s102 sub
|
|
r3 = r3 - r6; // s103 sub
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r6, 4u); r1 = r1 ^ t_; } // s104 shfl
|
|
r3 = r3 | r5; // s105 or
|
|
r0 = r0 + r1 + ((((sel >> 22u) & 1u) != 0u) ? 0x018722b4u : 0x9a16bcfbu); // s106 add
|
|
r2 = r2 + r5 + ((((sel >> 6u) & 1u) != 0u) ? 0x44cbb3d8u : 0xbc4b5e44u); // s107 add
|
|
r2 = r6 * r1 + r2; // s108 mad
|
|
r0 = r0 ^ r1; // s109 xor
|
|
r4 = r4 | r2; // s110 or
|
|
r2 = rotl_imm(r2, 13u); // s111 rotl
|
|
r5 = mul_hi(r5, r2); // s112 mulhi
|
|
r5 = r5 ^ r1; // s113 xor
|
|
r0 = rotl_imm(r0, 15u); // s114 rotl
|
|
r7 = r7 * r0; // s115 mul
|
|
r0 = r7 * r0 + r0; // s116 mad
|
|
r4 = rotl_imm(r4, 12u); // s117 rotl
|
|
r1 = r1 * r6; // s118 mul
|
|
r0 = mul_hi(r0, r3); // s119 mulhi
|
|
r4 = r0 * r0 + r4; // s120 mad
|
|
r0 = r0 + r5 + ((((sel >> 16u) & 1u) != 0u) ? 0x153e7b81u : 0x227a1887u); // s121 add
|
|
r6 = r6 + r3 + ((((sel >> 31u) & 1u) != 0u) ? 0x537e0843u : 0xfbb1908bu); // s122 add
|
|
r4 = mul_hi(r4, r5); // s123 mulhi
|
|
r3 = r3 ^ r1; // s124 xor
|
|
r3 = r3 + r7 + ((((sel >> 15u) & 1u) != 0u) ? 0xc2875857u : 0xc3e1337du); // s125 add
|
|
r4 = rotr_var(r4, r7); // s126 rotr
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r0, 2u); r1 = r1 ^ t_; } // s127 shfl
|
|
r3 = rotr_var(r3, r6); // s128 rotr
|
|
r1 = r1 * r0; // s129 mul
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r1, 16u); r4 = r4 ^ t_; } // s130 shfl
|
|
r4 = r4 + r0 + ((((sel >> 5u) & 1u) != 0u) ? 0x39900c5eu : 0x87bd1ad9u); // s131 add
|
|
r3 = r0 * r3 + r3; // s132 mad
|
|
r4 = r4 * r2; // s133 mul
|
|
r5 = r6 * r0 + r5; // s134 mad
|
|
r4 = r5 * r4 + r4; // s135 mad
|
|
r6 = r6 + r3 + ((((sel >> 28u) & 1u) != 0u) ? 0xcb7e81cau : 0xc59dd71du); // s136 add
|
|
r7 = r7 * r5; // s137 mul
|
|
r6 = r6 * r3; // s138 mul
|
|
r0 = rotr_var(r0, r4); // s139 rotr
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r5, 16u); r3 = r3 ^ t_; } // s140 shfl
|
|
r6 = rotr_var(r6, r7); // s141 rotr
|
|
r3 = r3 * r7; // s142 mul
|
|
r0 = r0 + r7 + ((((sel >> 9u) & 1u) != 0u) ? 0x602dc90du : 0x273f8ee2u); // s143 add
|
|
r5 = rotl_imm(r5, 26u); // s144 rotl
|
|
r2 = r2 ^ r4; // s145 xor
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r5, 16u); r6 = r6 ^ t_; } // s146 shfl
|
|
r1 = r1 * r4; // s147 mul
|
|
r2 = r1 * r7 + r2; // s148 mad
|
|
r7 = rotr_var(r7, r2); // s149 rotr
|
|
r7 = rotr_var(r7, r3); // s150 rotr
|
|
r5 = r5 + r2 + ((((sel >> 17u) & 1u) != 0u) ? 0xb3e8ca69u : 0x6f435830u); // s151 add
|
|
r6 = r6 ^ r2; // s152 xor
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r2, 16u); r1 = r1 ^ t_; } // s153 shfl
|
|
r2 = r2 ^ r6; // s154 xor
|
|
r5 = rotr_var(r5, r7); // s155 rotr
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r3, 4u); r1 = r1 ^ t_; } // s156 shfl
|
|
r6 = r6 ^ r5; // s157 xor
|
|
r0 = r0 ^ r7; // s158 xor
|
|
r0 = r0 ^ r7; // s159 xor
|
|
r0 = mul_hi(r0, r3); // s160 mulhi
|
|
r1 = r1 * r5; // s161 mul
|
|
r4 = rotr_var(r4, r0); // s162 rotr
|
|
r4 = r1 * r2 + r4; // s163 mad
|
|
r3 = r3 + r6 + ((((sel >> 18u) & 1u) != 0u) ? 0xe88bec07u : 0x7b5c68f7u); // s164 add
|
|
r0 = rotl_imm(r0, 13u); // s165 rotl
|
|
r2 = r2 ^ r6; // s166 xor
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r4, 16u); r5 = r5 ^ t_; } // s167 shfl
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r7, 2u); r2 = r2 ^ t_; } // s168 shfl
|
|
r7 = r7 ^ r6; // s169 xor
|
|
r0 = r7 * r2 + r0; // s170 mad
|
|
r3 = rotl_imm(r3, 3u); // s171 rotl
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r6, 2u); r7 = r7 ^ t_; } // s172 shfl
|
|
r0 = r0 + r1 + ((((sel >> 28u) & 1u) != 0u) ? 0xadc930fcu : 0xc53a1209u); // s173 add
|
|
r0 = r0 * r6; // s174 mul
|
|
r7 = mul_hi(r7, r0); // s175 mulhi
|
|
r3 = r3 | r2; // s176 or
|
|
r4 = r4 + r3 + ((((sel >> 16u) & 1u) != 0u) ? 0x36cdb68eu : 0x2def94b8u); // s177 add
|
|
r6 = r6 ^ r2; // s178 xor
|
|
r0 = rotl_imm(r0, 16u); // s179 rotl
|
|
r4 = r4 + r3 + ((((sel >> 5u) & 1u) != 0u) ? 0x230f4372u : 0x774065efu); // s180 add
|
|
r6 = r2 * r2 + r6; // s181 mad
|
|
r4 = r4 + r5 + ((((sel >> 18u) & 1u) != 0u) ? 0xbc379c7cu : 0x8c37e75bu); // s182 add
|
|
r0 = rotl_imm(r0, 26u); // s183 rotl
|
|
r4 = r4 | r2; // s184 or
|
|
r0 = r0 * r2; // s185 mul
|
|
r3 = r3 | r5; // s186 or
|
|
r1 = mul_hi(r1, r0); // s187 mulhi
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r2, 16u); r4 = r4 ^ t_; } // s188 shfl
|
|
r5 = rotr_var(r5, r4); // s189 rotr
|
|
r3 = r0 * r3 + r3; // s190 mad
|
|
r1 = r1 | r7; // s191 or
|
|
r7 = r7 | r1; // s192 or
|
|
r1 = r5 * r4 + r1; // s193 mad
|
|
r0 = r0 - r7; // s194 sub
|
|
r6 = r6 + r0 + ((((sel >> 9u) & 1u) != 0u) ? 0x8751e547u : 0x2f8a59eeu); // s195 add
|
|
r0 = rotl_imm(r0, 8u); // s196 rotl
|
|
r3 = r3 + r4 + ((((sel >> 2u) & 1u) != 0u) ? 0x02b9bb4cu : 0x0f369a86u); // s197 add
|
|
r4 = r4 + r2 + ((((sel >> 11u) & 1u) != 0u) ? 0x74240d4cu : 0xe7922061u); // s198 add
|
|
r2 = r2 * r5; // s199 mul
|
|
r6 = r6 + r7 + ((((sel >> 16u) & 1u) != 0u) ? 0x7649c3c3u : 0xee0e3356u); // s200 add
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r1, 16u); r6 = r6 ^ t_; } // s201 shfl
|
|
r4 = r4 * r2; // s202 mul
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r2, 1u); r3 = r3 ^ t_; } // s203 shfl
|
|
r7 = r2 * r1 + r7; // s204 mad
|
|
r2 = rotl_imm(r2, 5u); // s205 rotl
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r1, 8u); r7 = r7 ^ t_; } // s206 shfl
|
|
r7 = r7 * r2; // s207 mul
|
|
r0 = r0 * r3; // s208 mul
|
|
r1 = rotl_imm(r1, 7u); // s209 rotl
|
|
r5 = r5 + r3 + ((((sel >> 23u) & 1u) != 0u) ? 0x3d8f8187u : 0xc8db10a0u); // s210 add
|
|
r5 = r5 - r0; // s211 sub
|
|
r0 = rotr_var(r0, r7); // s212 rotr
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r2, 8u); r4 = r4 ^ t_; } // s213 shfl
|
|
r1 = r1 ^ r3; // s214 xor
|
|
r1 = rotl_imm(r1, 19u); // s215 rotl
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r3, 2u); r6 = r6 ^ t_; } // s216 shfl
|
|
r2 = mul_hi(r2, r5); // s217 mulhi
|
|
r5 = rotl_imm(r5, 14u); // s218 rotl
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r1, 8u); r2 = r2 ^ t_; } // s219 shfl
|
|
r7 = r7 - r5; // s220 sub
|
|
r3 = mul_hi(r3, r6); // s221 mulhi
|
|
r7 = r7 | r1; // s222 or
|
|
r1 = mul_hi(r1, r5); // s223 mulhi
|
|
r7 = r7 + r5 + ((((sel >> 24u) & 1u) != 0u) ? 0xd5a3fb54u : 0x689cdcf5u); // s224 add
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r7, 16u); r5 = r5 ^ t_; } // s225 shfl
|
|
r3 = mul_hi(r3, r2); // s226 mulhi
|
|
r0 = r0 ^ r2; // s227 xor
|
|
r7 = r7 + r4 + ((((sel >> 8u) & 1u) != 0u) ? 0xba3b9728u : 0x267f928du); // s228 add
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r7, 2u); r1 = r1 ^ t_; } // s229 shfl
|
|
r4 = r4 - r6; // s230 sub
|
|
r0 = r0 | r1; // s231 or
|
|
r2 = rotl_imm(r2, 10u); // s232 rotl
|
|
r4 = r4 * r7; // s233 mul
|
|
r6 = r0 * r0 + r6; // s234 mad
|
|
r4 = rotl_imm(r4, 29u); // s235 rotl
|
|
r7 = r7 + r2 + ((((sel >> 13u) & 1u) != 0u) ? 0xa437db0eu : 0xed6dd62au); // s236 add
|
|
r4 = rotr_var(r4, r7); // s237 rotr
|
|
r2 = r2 + r0 + ((((sel >> 17u) & 1u) != 0u) ? 0x1c000e82u : 0x9f612006u); // s238 add
|
|
r7 = mul_hi(r7, r5); // s239 mulhi
|
|
r3 = mul_hi(r3, r6); // s240 mulhi
|
|
r0 = r0 ^ r7; // s241 xor
|
|
r4 = rotl_imm(r4, 11u); // s242 rotl
|
|
r3 = rotl_imm(r3, 21u); // s243 rotl
|
|
r4 = r4 + r2 + ((((sel >> 13u) & 1u) != 0u) ? 0x12705678u : 0x83ba196cu); // s244 add
|
|
r7 = r7 + r5 + ((((sel >> 2u) & 1u) != 0u) ? 0xea712528u : 0xa5d39c13u); // s245 add
|
|
r0 = r0 * r5; // s246 mul
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r0, 2u); r4 = r4 ^ t_; } // s247 shfl
|
|
r3 = r3 ^ r2; // s248 xor
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r3, 4u); r7 = r7 ^ t_; } // s249 shfl
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r3, 16u); r5 = r5 ^ t_; } // s250 shfl
|
|
r5 = r5 + r3 + ((((sel >> 21u) & 1u) != 0u) ? 0x26ce965fu : 0x3006c6ebu); // s251 add
|
|
r3 = rotl_imm(r3, 1u); // s252 rotl
|
|
r7 = mul_hi(r7, r1); // s253 mulhi
|
|
r1 = r1 + r5 + ((((sel >> 1u) & 1u) != 0u) ? 0x9e34c13fu : 0xc2f46c6du); // s254 add
|
|
r4 = rotr_var(r4, r7); // s255 rotr
|
|
}
|
|
}
|
|
uint lo = r0 ^ rotl_imm(r1, 7u) ^ rotl_imm(r2, 14u) ^ rotl_imm(r3, 21u);
|
|
uint hi = r4 ^ rotl_imm(r5, 9u) ^ rotl_imm(r6, 18u) ^ rotl_imm(r7, 27u);
|
|
out[gid] = ((ulong)hi << 32) | (ulong)lo;
|
|
}
|
|
|
|
#if IGNEUM_EXCHANGE != 0
|
|
// Reports the sub-group size this device uses for a work-group of IGNEUM_GROUP items. host.c runs it only when the
|
|
// per-kernel query (clGetKernelSubGroupInfoKHR on igneum_hash) is unavailable; that query is preferred because a
|
|
// compiler may pick a different wave width per kernel (RDNA: wave32 or wave64). See WAVEFRONT.md.
|
|
IGNEUM_KERNEL_HASH void igneum_probe_subgroup(__global uint* out) {
|
|
if (get_local_id(0) == 0u) { out[0] = get_sub_group_size(); out[1] = get_num_sub_groups(); }
|
|
}
|
|
#endif
|