771e89abce
Nothing changes for the default class: the pinned packs are byte-identical (tests/packs.rs), the v2 draw stream is untouched.
LoadClass {mix, load_slots, scratch}: fixed widths w16, w64, w64x4 (4 loads of 64 B), era mixes 50/35/15 and 25/50/25 drawn per load with one extra below(100) roll, and the scratch variant scr0/2/4/8 (persistent warps, 1 MiB per warp, tagged lazy fill, measurement only). A wide load reads the W-aligned address and folds every word: x = dst ^ w0; x = (rotl(x, 11) * 0x9e3779b1) ^ w[j]. Program ids carry the class. proto-opencl/host.c taken from opencl-rdna4 5034e81 (--memprobe, select read-back).
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
278 lines
18 KiB
Common Lisp
278 lines
18 KiB
Common Lisp
// Generated by igneum-pow export (generator v2) for seed "igneum-genesis". 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 seed-parameterised mixer + one 64-byte cache read, then a final mixer. 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] = 0x3067619fu ^ prev[4];
|
|
x[5] = 0x3c269176u ^ prev[5];
|
|
x[6] = 0x84a03b03u ^ prev[6];
|
|
x[7] = 0xf8c63294u ^ prev[7];
|
|
x[8] = 0xff977c5bu ^ prev[8];
|
|
x[9] = 0xe60def3eu ^ prev[9];
|
|
x[10] = 0x63630141u ^ prev[10];
|
|
x[11] = 0xb8fbcb58u ^ 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] ^ (0xbab68293u + rk)) * 0x42146205u;
|
|
s[1] = (s[1] ^ (0xcc162340u + rk)) * 0x52cbe0fbu;
|
|
s[2] = (s[2] ^ (0x6ce151ccu + rk)) * 0x7ecf4a03u;
|
|
s[3] = (s[3] ^ (0xe62b8997u + rk)) * 0x6728907fu;
|
|
s[4] = (s[4] ^ (0xc9c80297u + rk)) * 0xd81d9751u;
|
|
s[5] = (s[5] ^ (0xf74a1654u + rk)) * 0x132952c3u;
|
|
s[6] = (s[6] ^ (0x3d704af5u + rk)) * 0xf60de277u;
|
|
s[7] = (s[7] ^ (0x3cf522b7u + rk)) * 0x05358035u;
|
|
s[8] = (s[8] ^ (0x2b9cac04u + rk)) * 0xbaf6499du;
|
|
s[9] = (s[9] ^ (0xa880ac10u + rk)) * 0xe4db9667u;
|
|
s[10] = (s[10] ^ (0x13e5dd1du + rk)) * 0x3e98f45du;
|
|
s[11] = (s[11] ^ (0x6fc3e233u + rk)) * 0xd0004eddu;
|
|
s[12] = (s[12] ^ (0x2d83eeacu + rk)) * 0x2691630du;
|
|
s[13] = (s[13] ^ (0x9006e8bfu + rk)) * 0x9beb3bcfu;
|
|
s[14] = (s[14] ^ (0x2c4b5362u + rk)) * 0xab310379u;
|
|
s[15] = (s[15] ^ (0x31b49ee2u + rk)) * 0x99cfb423u;
|
|
MH_QR(s[0], s[4], s[8], s[12], 20u, 20u, 19u, 4u) MH_QR(s[1], s[5], s[9], s[13], 20u, 20u, 19u, 4u)
|
|
MH_QR(s[2], s[6], s[10], s[14], 20u, 20u, 19u, 4u) MH_QR(s[3], s[7], s[11], s[15], 20u, 20u, 19u, 4u)
|
|
MH_QR(s[0], s[5], s[10], s[15], 26u, 3u, 3u, 27u) MH_QR(s[1], s[6], s[11], s[12], 26u, 3u, 3u, 27u)
|
|
MH_QR(s[2], s[7], s[8], s[13], 26u, 3u, 3u, 27u) MH_QR(s[3], s[4], s[9], s[14], 26u, 3u, 3u, 27u)
|
|
}
|
|
|
|
// Item t: 16 words. s = (K, t * MUL[i] + RC[i]); 8 rounds of mixer + cache line s[0] & mask; final mixer.
|
|
static inline void mh_item(__global const uint* cache, uint t, uint* s) {
|
|
s[0] = 0x3067619fu;
|
|
s[1] = 0x3c269176u;
|
|
s[2] = 0x84a03b03u;
|
|
s[3] = 0xf8c63294u;
|
|
s[4] = 0xff977c5bu;
|
|
s[5] = 0xe60def3eu;
|
|
s[6] = 0x63630141u;
|
|
s[7] = 0xb8fbcb58u;
|
|
s[8] = t * 0x42146205u + 0xbab68293u;
|
|
s[9] = t * 0x52cbe0fbu + 0xcc162340u;
|
|
s[10] = t * 0x7ecf4a03u + 0x6ce151ccu;
|
|
s[11] = t * 0x6728907fu + 0xe62b8997u;
|
|
s[12] = t * 0xd81d9751u + 0xc9c80297u;
|
|
s[13] = t * 0x132952c3u + 0xf74a1654u;
|
|
s[14] = t * 0xf60de277u + 0x3d704af5u;
|
|
s[15] = t * 0x05358035u + 0x3cf522b7u;
|
|
for (uint r = 0u; r < 8u; ++r) {
|
|
mh_mixer(s, 0x9E3779B9u * (r + 1u));
|
|
__global const uint* line = cache + ((s[0] & MH_CACHE_LINE_MASK) * 16u);
|
|
for (uint i = 0u; i < 16u; ++i) s[i] ^= line[i];
|
|
}
|
|
mh_mixer(s, 0x9E3779B9u * 9u);
|
|
}
|
|
// dataset[w] without the dataset: derive item w >> 4 and take word w & 15.
|
|
static inline uint mh_word(__global const uint* cache, uint w) { uint s[16]; mh_item(cache, w >> 4u, s); return s[w & 15u]; }
|
|
|
|
// 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);
|
|
__global uint* d = ds + ((ulong)t * 16u);
|
|
for (uint i = 0u; i < 16u; ++i) d[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 ^ 0x67a9a7beu; x += 0x9e3779b9u; x = splitmix32(x); r0 = x ^ 0x1a155b25u; } // SEEDW[0], 0x9e3779b9u * 1u, SEEDW[1]
|
|
{ uint x = nonce ^ 0x1a155b25u; x += 0x3c6ef372u; x = splitmix32(x); r1 = x ^ 0xfddfb732u; } // SEEDW[1], 0x9e3779b9u * 2u, SEEDW[2]
|
|
{ uint x = nonce ^ 0xfddfb732u; x += 0xdaa66d2bu; x = splitmix32(x); r2 = x ^ 0x4b5af2e8u; } // SEEDW[2], 0x9e3779b9u * 3u, SEEDW[3]
|
|
{ uint x = nonce ^ 0x4b5af2e8u; x += 0x78dde6e4u; x = splitmix32(x); r3 = x ^ 0xc55caf33u; } // SEEDW[3], 0x9e3779b9u * 4u, SEEDW[4]
|
|
{ uint x = nonce ^ 0xc55caf33u; x += 0x1715609du; x = splitmix32(x); r4 = x ^ 0xa27c13b7u; } // SEEDW[4], 0x9e3779b9u * 5u, SEEDW[5]
|
|
{ uint x = nonce ^ 0xa27c13b7u; x += 0xb54cda56u; x = splitmix32(x); r5 = x ^ 0x06628a48u; } // SEEDW[5], 0x9e3779b9u * 6u, SEEDW[6]
|
|
{ uint x = nonce ^ 0x06628a48u; x += 0x5384540fu; x = splitmix32(x); r6 = x ^ 0x03852469u; } // SEEDW[6], 0x9e3779b9u * 7u, SEEDW[7]
|
|
{ uint x = nonce ^ 0x03852469u; x += 0xf1bbcdc8u; x = splitmix32(x); r7 = x ^ 0x67a9a7beu; } // SEEDW[7], 0x9e3779b9u * 8u, SEEDW[0]
|
|
|
|
for (uint it = 0u; it < 8u; ++it) {
|
|
uint sel = r0;
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r6, 4u); r0 = r0 ^ t_; } // 0 shfl
|
|
r7 = mul_hi(r7, r6); // 1 mulhi
|
|
r5 = r5 + r7 + ((((sel >> 21u) & 1u) != 0u) ? 0xdd04a5dau : 0xe9239829u); // 2 add
|
|
r2 = r3 * r2 + r2; // 3 mad
|
|
r0 = mul_hi(r0, r7); // 4 mulhi
|
|
r5 = r5 + r7 + ((((sel >> 7u) & 1u) != 0u) ? 0x9043323eu : 0x265677dcu); // 5 add
|
|
r6 = rotr_var(r6, r5); // 6 rotr
|
|
r1 = r1 * r4; // 7 mul
|
|
r4 = r4 ^ r0; // 8 xor
|
|
r5 = r5 + r2 + ((((sel >> 7u) & 1u) != 0u) ? 0x504c2692u : 0xc0aca976u); // 9 add
|
|
r7 = mul_hi(r7, r5); // 10 mulhi
|
|
r4 = r4 ^ r2; // 11 xor
|
|
r0 = mul_hi(r0, r4); // 12 mulhi
|
|
r2 = mul_hi(r2, r5); // 13 mulhi
|
|
r3 = mul_hi(r3, r0); // 14 mulhi
|
|
r1 = rotl_imm(r1, 22u); // 15 rotl
|
|
{ uint b_ = (r4 & mask) & ~15u; uint4 v0_ = vload4(0u, ds + b_); uint4 v1_ = vload4(1u, ds + b_); uint4 v2_ = vload4(2u, ds + b_); uint4 v3_ = vload4(3u, ds + b_); uint x_ = r6 ^ v0_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.w; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.w; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.w; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.w; r6 = x_; } // 16 load
|
|
r0 = rotl_imm(r0, 11u); // 17 rotl
|
|
r6 = r6 + r1 + ((((sel >> 7u) & 1u) != 0u) ? 0x0d110a7au : 0xa377d905u); // 18 add
|
|
r4 = r4 + r6 + ((((sel >> 2u) & 1u) != 0u) ? 0x9b05eed7u : 0x56170e75u); // 19 add
|
|
r2 = r2 + r4 + ((((sel >> 27u) & 1u) != 0u) ? 0x3ac915d2u : 0xdc8d29d5u); // 20 add
|
|
r6 = mul_hi(r6, r2); // 21 mulhi
|
|
r6 = r6 - r2; // 22 sub
|
|
r7 = rotr_var(r7, r1); // 23 rotr
|
|
r0 = rotr_var(r0, r2); // 24 rotr
|
|
r3 = r2 * r7 + r3; // 25 mad
|
|
{ uint t_; IGNEUM_SHFL_XOR(t_, r6, 1u); r3 = r3 ^ t_; } // 26 shfl
|
|
r5 = rotl_imm(r5, 15u); // 27 rotl
|
|
r6 = r6 * r4; // 28 mul
|
|
r2 = r2 + r0 + ((((sel >> 14u) & 1u) != 0u) ? 0x09ed045eu : 0x2d1d020au); // 29 add
|
|
r1 = r1 - r6; // 30 sub
|
|
r0 = mul_hi(r0, r5); // 31 mulhi
|
|
{ uint b_ = (r0 & mask) & ~15u; uint4 v0_ = vload4(0u, ds + b_); uint4 v1_ = vload4(1u, ds + b_); uint4 v2_ = vload4(2u, ds + b_); uint4 v3_ = vload4(3u, ds + b_); uint x_ = r3 ^ v0_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.w; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.w; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.w; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.w; r3 = x_; } // 32 load
|
|
r7 = r7 ^ r5; // 33 xor
|
|
{ uint b_ = (r4 & mask) & ~15u; uint4 v0_ = vload4(0u, ds + b_); uint4 v1_ = vload4(1u, ds + b_); uint4 v2_ = vload4(2u, ds + b_); uint4 v3_ = vload4(3u, ds + b_); uint x_ = r5 ^ v0_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.w; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.w; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.w; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.w; r5 = x_; } // 34 load
|
|
r4 = rotr_var(r4, r0); // 35 rotr
|
|
r3 = r3 + r0 + ((((sel >> 21u) & 1u) != 0u) ? 0xd1b7c4e1u : 0x4c115f69u); // 36 add
|
|
r4 = r4 + r3 + ((((sel >> 19u) & 1u) != 0u) ? 0xaca40f85u : 0x4e8977ceu); // 37 add
|
|
r4 = r1 * r4 + r4; // 38 mad
|
|
r6 = r6 - r7; // 39 sub
|
|
r2 = r2 * r0; // 40 mul
|
|
r7 = r7 * r2; // 41 mul
|
|
r3 = r3 + r0 + ((((sel >> 9u) & 1u) != 0u) ? 0x87d9ef84u : 0x028aa63cu); // 42 add
|
|
r5 = r5 ^ r3; // 43 xor
|
|
r6 = rotl_imm(r6, 20u); // 44 rotl
|
|
r5 = r5 + r3 + ((((sel >> 31u) & 1u) != 0u) ? 0xff9a2d77u : 0x89948c4bu); // 45 add
|
|
r2 = r5 * r1 + r2; // 46 mad
|
|
r0 = r0 + r1 + ((((sel >> 6u) & 1u) != 0u) ? 0xa8848b30u : 0x0be29eaau); // 47 add
|
|
r0 = r0 + r2 + ((((sel >> 6u) & 1u) != 0u) ? 0x699fd448u : 0x4f92b968u); // 48 add
|
|
r3 = r3 + r5 + ((((sel >> 6u) & 1u) != 0u) ? 0xb17aad78u : 0x77b1520du); // 49 add
|
|
r0 = r6 * r0 + r0; // 50 mad
|
|
r2 = r2 + r0 + ((((sel >> 22u) & 1u) != 0u) ? 0x5c933bd0u : 0x14bde8e1u); // 51 add
|
|
r7 = r7 - r0; // 52 sub
|
|
r7 = r7 + r4 + ((((sel >> 11u) & 1u) != 0u) ? 0x546f4095u : 0x7149e3c9u); // 53 add
|
|
r2 = r2 * r5; // 54 mul
|
|
r6 = r6 ^ r2; // 55 xor
|
|
r4 = r4 ^ r3; // 56 xor
|
|
r4 = r4 + r6 + ((((sel >> 27u) & 1u) != 0u) ? 0x1e07c3d9u : 0x89841d87u); // 57 add
|
|
r0 = r0 | r6; // 58 or
|
|
{ uint b_ = (r4 & mask) & ~15u; uint4 v0_ = vload4(0u, ds + b_); uint4 v1_ = vload4(1u, ds + b_); uint4 v2_ = vload4(2u, ds + b_); uint4 v3_ = vload4(3u, ds + b_); uint x_ = r1 ^ v0_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v0_.w; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v1_.w; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v2_.w; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.x; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.y; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.z; x_ = (rotl_imm(x_, 11u) * 0x9e3779b1u) ^ v3_.w; r1 = x_; } // 59 load
|
|
r5 = r5 ^ r2; // 60 xor
|
|
r0 = r0 + r3 + ((((sel >> 22u) & 1u) != 0u) ? 0xe686f917u : 0x83245ee9u); // 61 add
|
|
r1 = mul_hi(r1, r4); // 62 mulhi
|
|
r7 = rotl_imm(r7, 12u); // 63 rotl
|
|
}
|
|
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
|