igneum/proto-cuda/packs/igneum-hourly/kernel.cl

// Generated by proto-metal/igneum-bench --export-pack for seed "igneum-hourly". 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;
}

// dataset[i] = ds_elem(i, d0, d1) for i < n. Same closed form as the Metal igneum_fill kernel.
__kernel void igneum_fill(__global uint* ds, uint n, uint d0, uint d1) {
    uint i = (uint)get_global_id(0);
    if (i < n) ds[i] = ds_elem(i, d0, d1);
}

// 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 ^ 0x6bdee811u; x += 0x9e3779b9u; x = splitmix32(x); r0 = x ^ 0x8f488bbeu; }  // SEEDW[0], 0x9e3779b9u * 1u, SEEDW[1]
    { uint x = nonce ^ 0x8f488bbeu; x += 0x3c6ef372u; x = splitmix32(x); r1 = x ^ 0xc5cdece7u; }  // SEEDW[1], 0x9e3779b9u * 2u, SEEDW[2]
    { uint x = nonce ^ 0xc5cdece7u; x += 0xdaa66d2bu; x = splitmix32(x); r2 = x ^ 0x210af22du; }  // SEEDW[2], 0x9e3779b9u * 3u, SEEDW[3]
    { uint x = nonce ^ 0x210af22du; x += 0x78dde6e4u; x = splitmix32(x); r3 = x ^ 0x2f687b65u; }  // SEEDW[3], 0x9e3779b9u * 4u, SEEDW[4]
    { uint x = nonce ^ 0x2f687b65u; x += 0x1715609du; x = splitmix32(x); r4 = x ^ 0x17471eeeu; }  // SEEDW[4], 0x9e3779b9u * 5u, SEEDW[5]
    { uint x = nonce ^ 0x17471eeeu; x += 0xb54cda56u; x = splitmix32(x); r5 = x ^ 0xee16e284u; }  // SEEDW[5], 0x9e3779b9u * 6u, SEEDW[6]
    { uint x = nonce ^ 0xee16e284u; x += 0x5384540fu; x = splitmix32(x); r6 = x ^ 0xfc9eb8f9u; }  // SEEDW[6], 0x9e3779b9u * 7u, SEEDW[7]
    { uint x = nonce ^ 0xfc9eb8f9u; x += 0xf1bbcdc8u; x = splitmix32(x); r7 = x ^ 0x6bdee811u; }  // SEEDW[7], 0x9e3779b9u * 8u, SEEDW[0]

    for (uint it = 0u; it < 8u; ++it) {
        uint sel = r0;
        r6 = r6 ^ ds[r5 & mask];  // 0 load
        r3 = r3 ^ r7;  // 1 xor
        r6 = mul_hi(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
        { uint t_; IGNEUM_SHFL_XOR(t_, r3, 8u); r0 = r0 ^ t_; }  // 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 = mul_hi(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 = mul_hi(r3, r4);  // 24 mulhi
        r6 = r6 ^ r1;  // 25 xor
        r1 = rotl_imm(r1, 31u);  // 26 rotl
        { uint t_; IGNEUM_SHFL_XOR(t_, r6, 4u); r3 = r3 ^ t_; }  // 27 shfl
        r6 = r6 - r5;  // 28 sub
        r6 = rotr_var(r6, r3);  // 29 rotr
        { uint t_; IGNEUM_SHFL_XOR(t_, r4, 4u); r0 = r0 ^ t_; }  // 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
        { uint t_; IGNEUM_SHFL_XOR(t_, r5, 8u); r4 = r4 ^ t_; }  // 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
        { uint t_; IGNEUM_SHFL_XOR(t_, r4, 16u); r5 = r5 ^ t_; }  // 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 = mul_hi(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 = mul_hi(r1, r5);  // 62 mulhi
        r7 = r7 + r3 + ((((sel >> 24u) & 1u) != 0u) ? 0x05e4fc1du : 0xc23e27c9u);  // 63 add
    }
    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