dc84789e34
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 23810df (--memprobe, select read-back).
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
109 lines
8.9 KiB
Metal
109 lines
8.9 KiB
Metal
#include <metal_stdlib>
|
|
using namespace metal;
|
|
|
|
#define MASK 0x0fffffffu
|
|
constant uint SEEDW[8] = { 0x67a9a7beu, 0x1a155b25u, 0xfddfb732u, 0x4b5af2e8u, 0xc55caf33u, 0xa27c13b7u, 0x06628a48u, 0x03852469u };
|
|
|
|
inline uint splitmix32(uint x) {
|
|
x ^= x >> 16; x *= 0x7feb352du;
|
|
x ^= x >> 15; x *= 0x846ca68bu;
|
|
x ^= x >> 16;
|
|
return x;
|
|
}
|
|
inline uint rotl_imm(uint x, uint n) { return (x << n) | (x >> (32u - n)); } // n in 1..31
|
|
inline uint rotr_var(uint x, uint n) { n &= 31u; return (x >> n) | (x << ((32u - n) & 31u)); }
|
|
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;
|
|
}
|
|
|
|
kernel void igneum_hash(device const uint* dataset [[buffer(0)]],
|
|
device ulong* out [[buffer(1)]],
|
|
constant uint& baseNonce [[buffer(2)]],
|
|
uint gid [[thread_position_in_grid]]) {
|
|
uint nonce = baseNonce + gid;
|
|
uint r0, r1, r2, r3, r4, r5, r6, r7;
|
|
{ uint x = nonce ^ SEEDW[0]; x += 0x9e3779b9u * 1u; x = splitmix32(x); r0 = x ^ SEEDW[1]; }
|
|
{ uint x = nonce ^ SEEDW[1]; x += 0x9e3779b9u * 2u; x = splitmix32(x); r1 = x ^ SEEDW[2]; }
|
|
{ uint x = nonce ^ SEEDW[2]; x += 0x9e3779b9u * 3u; x = splitmix32(x); r2 = x ^ SEEDW[3]; }
|
|
{ uint x = nonce ^ SEEDW[3]; x += 0x9e3779b9u * 4u; x = splitmix32(x); r3 = x ^ SEEDW[4]; }
|
|
{ uint x = nonce ^ SEEDW[4]; x += 0x9e3779b9u * 5u; x = splitmix32(x); r4 = x ^ SEEDW[5]; }
|
|
{ uint x = nonce ^ SEEDW[5]; x += 0x9e3779b9u * 6u; x = splitmix32(x); r5 = x ^ SEEDW[6]; }
|
|
{ uint x = nonce ^ SEEDW[6]; x += 0x9e3779b9u * 7u; x = splitmix32(x); r6 = x ^ SEEDW[7]; }
|
|
{ uint x = nonce ^ SEEDW[7]; x += 0x9e3779b9u * 8u; x = splitmix32(x); r7 = x ^ SEEDW[0]; }
|
|
|
|
for (uint it = 0u; it < 8u; ++it) {
|
|
uint sel = r0;
|
|
r2 = r3 * r4 + r2; // 0
|
|
r1 = r1 + r7 + select(0x42da7657u, 0xc3bd2355u, ((sel >> 4u) & 1u) != 0u); // 1
|
|
r2 = r2 + r3 + select(0x61f0b51cu, 0x2735a174u, ((sel >> 26u) & 1u) != 0u); // 2
|
|
r4 = r0 * r6 + r4; // 3
|
|
{ uint b_ = (r2 & MASK) & ~3u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; uint x_ = r7 ^ 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; r7 = x_; } // 4
|
|
{ uint b_ = (r1 & MASK) & ~3u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; uint x_ = r4 ^ 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; r4 = x_; } // 5
|
|
r6 = r6 ^ simd_shuffle_xor(r3, (ushort)4); // 6
|
|
r1 = r1 ^ simd_shuffle_xor(r5, (ushort)8); // 7
|
|
r7 = r7 ^ r5; // 8
|
|
r3 = r3 | r4; // 9
|
|
r1 = r1 | r2; // 10
|
|
{ uint b_ = (r3 & MASK) & ~3u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; uint x_ = r4 ^ 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; r4 = x_; } // 11
|
|
r6 = r6 | r2; // 12
|
|
r2 = r2 * r5; // 13
|
|
{ uint b_ = (r2 & MASK) & ~3u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; 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; r1 = x_; } // 14
|
|
r7 = rotl_imm(r7, 1u); // 15
|
|
{ uint b_ = (r6 & MASK) & ~3u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; 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; r3 = x_; } // 16
|
|
{ uint b_ = (r4 & MASK) & ~3u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; uint x_ = r7 ^ 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; r7 = x_; } // 17
|
|
r3 = r3 ^ simd_shuffle_xor(r4, (ushort)2); // 18
|
|
r4 = r0 * r2 + r4; // 19
|
|
r0 = r0 ^ simd_shuffle_xor(r6, (ushort)8); // 20
|
|
r5 = r5 ^ r7; // 21
|
|
r2 = mulhi(r2, r5); // 22
|
|
{ uint b_ = (r7 & MASK) & ~3u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; 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; r3 = x_; } // 23
|
|
r7 = mulhi(r7, r3); // 24
|
|
r5 = r5 | r4; // 25
|
|
r4 = r5 * r2 + r4; // 26
|
|
r5 = r5 * r1; // 27
|
|
r6 = mulhi(r6, r7); // 28
|
|
r6 = r6 + r1 + select(0x3b2d2124u, 0x187a9128u, ((sel >> 9u) & 1u) != 0u); // 29
|
|
r6 = rotr_var(r6, r7); // 30
|
|
{ uint b_ = (r1 & MASK) & ~3u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; 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; r3 = x_; } // 31
|
|
{ uint b_ = (r0 & MASK) & ~3u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; 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; r1 = x_; } // 32
|
|
r0 = r0 + r4 + select(0x2c35699fu, 0x351dde38u, ((sel >> 18u) & 1u) != 0u); // 33
|
|
{ uint b_ = (r2 & MASK) & ~3u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; uint x_ = r0 ^ 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; r0 = x_; } // 34
|
|
r0 = r0 * r3; // 35
|
|
r2 = r2 ^ r5; // 36
|
|
{ uint b_ = (r0 & MASK) & ~3u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; uint x_ = r4 ^ 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; r4 = x_; } // 37
|
|
r1 = r3 * r5 + r1; // 38
|
|
r0 = r0 + r3 + select(0xa907b90bu, 0x1b053acfu, ((sel >> 25u) & 1u) != 0u); // 39
|
|
r2 = rotr_var(r2, r5); // 40
|
|
r3 = r3 * r2; // 41
|
|
r1 = r1 + r5 + select(0xa32e000cu, 0x6058c2e3u, ((sel >> 20u) & 1u) != 0u); // 42
|
|
r3 = r3 ^ r4; // 43
|
|
{ uint b_ = (r5 & MASK) & ~3u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; 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; r3 = x_; } // 44
|
|
r1 = r1 + r5 + select(0x81b8bc2cu, 0x1907970cu, ((sel >> 7u) & 1u) != 0u); // 45
|
|
r7 = r7 ^ r1; // 46
|
|
r0 = r0 + r3 + select(0x838b5065u, 0x36360066u, ((sel >> 31u) & 1u) != 0u); // 47
|
|
r7 = mulhi(r7, r5); // 48
|
|
{ uint b_ = (r2 & MASK) & ~3u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; uint x_ = r0 ^ 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; r0 = x_; } // 49
|
|
r2 = r2 - r6; // 50
|
|
r7 = r7 - r5; // 51
|
|
r2 = r2 ^ r3; // 52
|
|
r7 = r7 - r0; // 53
|
|
r3 = r5 * r0 + r3; // 54
|
|
r7 = r7 ^ r5; // 55
|
|
{ uint b_ = (r7 & MASK) & ~3u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; uint x_ = r2 ^ 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; r2 = x_; } // 56
|
|
r5 = r5 - r6; // 57
|
|
{ uint b_ = (r3 & MASK) & ~3u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; 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; r1 = x_; } // 58
|
|
{ uint b_ = (r4 & MASK) & ~3u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; 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; r1 = x_; } // 59
|
|
r4 = r4 - r6; // 60
|
|
r1 = r1 * r2; // 61
|
|
r3 = r6 * r0 + r3; // 62
|
|
r0 = r0 + r1 + select(0x2fe0e98bu, 0xc88e2942u, ((sel >> 16u) & 1u) != 0u); // 63
|
|
}
|
|
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;
|
|
}
|