igneum-network/igneum
1
0
Fork 0
Code Issues Pull requests Releases Activity
igneum/proto-cuda/packs-readwidth/mixB-2/program.metal
igneum-labs dc84789e34 read-width experiment (gate 1): load classes W=4/16/64, per-load width mix, scratch RMW variant behind a generator flag; 20 packs; CPU verifier and acceptance mirror; emulator shims 
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>
2026-10-05 19:46:54 +00:00

109 lines
14 KiB
Metal
Raw Permalink Blame History

#include <metal_stdlib>
using namespace metal;
#define MASK 0x0fffffffu
constant uint SEEDW[8] = { 0x8c11f8d8u, 0x945e59f8u, 0x4ab18392u, 0xa6381ed7u, 0x797fadefu, 0x188a03acu, 0xea7032e0u, 0xcaf097f4u };
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;
r0 = r0 ^ r6; // 0
r0 = r1 * r4 + r0; // 1
r2 = rotl_imm(r2, 21u); // 2
r3 = rotr_var(r3, r0); // 3
r0 = r0 + r2 + select(0x7a34b651u, 0x474ae392u, ((sel >> 7u) & 1u) != 0u); // 4
r0 = mulhi(r0, r3); // 5
r3 = r3 ^ simd_shuffle_xor(r6, (ushort)4); // 6
r7 = rotr_var(r7, r2); // 7
{ uint b_ = (r0 & MASK) & ~15u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; uint4 v1_ = l_[1]; uint4 v2_ = l_[2]; uint4 v3_ = l_[3]; 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; 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; r2 = x_; } // 8
r4 = r4 | r5; // 9
{ uint b_ = (r4 & MASK) & ~15u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; uint4 v1_ = l_[1]; uint4 v2_ = l_[2]; uint4 v3_ = l_[3]; 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; 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; r0 = x_; } // 10
{ uint b_ = (r0 & MASK) & ~3u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; 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; r5 = x_; } // 11
{ uint b_ = (r7 & 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_; } // 12
r4 = r4 + r0 + select(0x99426786u, 0xf58f0389u, ((sel >> 16u) & 1u) != 0u); // 13
r0 = r0 + r6 + select(0x6c662222u, 0xb5bbd552u, ((sel >> 31u) & 1u) != 0u); // 14
r4 = rotr_var(r4, r6); // 15
r4 = r4 ^ r1; // 16
r3 = rotl_imm(r3, 22u); // 17
{ uint b_ = (r4 & 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_; } // 18
r3 = rotr_var(r3, r0); // 19
{ 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_; } // 20
r4 = rotl_imm(r4, 8u); // 21
{ uint b_ = (r5 & MASK) & ~15u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; uint4 v1_ = l_[1]; uint4 v2_ = l_[2]; uint4 v3_ = l_[3]; 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_; } // 22
r0 = rotr_var(r0, r3); // 23
r4 = r4 + r1 + select(0x05f4dd31u, 0xd1ff184bu, ((sel >> 13u) & 1u) != 0u); // 24
r4 = r4 ^ r2; // 25
r3 = mulhi(r3, r2); // 26
r5 = r5 - r7; // 27
{ uint b_ = (r5 & MASK) & ~15u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; uint4 v1_ = l_[1]; uint4 v2_ = l_[2]; uint4 v3_ = l_[3]; 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_; } // 28
r5 = r5 | r7; // 29
r4 = rotl_imm(r4, 28u); // 30
r3 = r3 ^ r5; // 31
r2 = r2 * r1; // 32
{ uint b_ = (r4 & 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_; } // 33
r5 = r5 ^ simd_shuffle_xor(r4, (ushort)16); // 34
{ uint b_ = (r5 & MASK) & ~3u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; 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; r6 = x_; } // 35
{ uint b_ = (r6 & MASK) & ~15u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; uint4 v1_ = l_[1]; uint4 v2_ = l_[2]; uint4 v3_ = l_[3]; 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; 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; r4 = x_; } // 36
r6 = r2 * r2 + r6; // 37
r6 = r6 ^ r3; // 38
r2 = r2 ^ dataset[r0 & MASK]; // 39
r1 = r1 * r2; // 40
{ uint b_ = (r2 & MASK) & ~15u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; uint4 v1_ = l_[1]; uint4 v2_ = l_[2]; uint4 v3_ = l_[3]; 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; 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; r4 = x_; } // 41
r4 = rotl_imm(r4, 31u); // 42
{ uint b_ = (r4 & MASK) & ~3u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; 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; r6 = x_; } // 43
r4 = r4 + r2 + select(0xb045a668u, 0xdd1539e6u, ((sel >> 24u) & 1u) != 0u); // 44
{ uint b_ = (r1 & MASK) & ~15u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; uint4 v1_ = l_[1]; uint4 v2_ = l_[2]; uint4 v3_ = l_[3]; 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_; } // 45
r7 = r7 + r6 + select(0x85c0036eu, 0x8a264d7du, ((sel >> 9u) & 1u) != 0u); // 46
r0 = r1 * r3 + r0; // 47
r7 = r7 | r0; // 48
r6 = r6 ^ r4; // 49
r1 = r1 + r7 + select(0x376c63cdu, 0x9d56aa91u, ((sel >> 11u) & 1u) != 0u); // 50
r5 = r5 * r4; // 51
r7 = rotr_var(r7, r0); // 52
r2 = r2 ^ r6; // 53
r0 = rotl_imm(r0, 22u); // 54
r3 = rotr_var(r3, r5); // 55
r6 = r6 + r2 + select(0x4face1f0u, 0x7fbd3c70u, ((sel >> 21u) & 1u) != 0u); // 56
r0 = r0 ^ simd_shuffle_xor(r1, (ushort)2); // 57
r2 = r6 * r0 + r2; // 58
r2 = mulhi(r2, r5); // 59
r5 = rotl_imm(r5, 2u); // 60
r6 = rotl_imm(r6, 28u); // 61
{ uint b_ = (r7 & MASK) & ~15u; device const uint4* l_ = (device const uint4*)(dataset + b_); uint4 v0_ = l_[0]; uint4 v1_ = l_[1]; uint4 v2_ = l_[2]; uint4 v3_ = l_[3]; 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; 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; r2 = x_; } // 62
r7 = mulhi(r7, r1); // 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;
}
Reference in a new issue View git blame Copy permalink