igneum-network/igneum
1
0
Fork 0
Code Issues Pull requests Releases Activity

read-width: OpenCL scratch kernels declare the __local exchange buffer before the unit loop (AMD's compiler requires the outermost scope; round 3 on the 9070 XT); seven scratch packs re-exported, vectors unchanged

Browse source 
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
This commit is contained in:
igneum-labs 2026-10-05 20:18:35 +00:00
parent a094972548
commit 357ccb9507
15 changed files with 194 additions and 157 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

85
igneum-pow/src/emit.rs
View file

@ -157,6 +157,14 @@ fn scratch_stmt(dialect: CoreDialect, d: &str, a: &str, slot_mask: u32) -> Strin
/// lottery hash's text is unchanged: `gid` is the unit's first output index plus the lane. The host MUST launch
/// `groups` as a multiple of N (a uniform trip count: the OpenCL local-memory exchange carries a barrier).
fn persistent_prologue(dialect: CoreDialect, words_per_lane: usize) -> String {
let (a, b) = persistent_prologue_parts(dialect, words_per_lane);
a + &b
}
/// The prologue in two parts: the warp's identity and arena, then the unit loop. OpenCL C requires a `__local`
/// variable at the outermost scope of the kernel (AMD's compiler enforces it, 5 October 2026, round 3 on the
/// 9070 XT), so the OpenCL kernels declare the exchange buffer between the two parts.
fn persistent_prologue_parts(dialect: CoreDialect, words_per_lane: usize) -> (String, String) {
let (u, tid, nthreads, ptr) = match dialect {
CoreDialect::Metal => ("uint", "tid", "nthreads", "device uint*"),
CoreDialect::Cuda => ("uint32_t", "(blockIdx.x * blockDim.x + threadIdx.x)", "(gridDim.x * blockDim.x)", "uint32_t*"),
@ -167,11 +175,12 @@ fn persistent_prologue(dialect: CoreDialect, words_per_lane: usize) -> String {
s.push_str(&format!(" {u} warp_ = {tid} >> 5;\n"));
s.push_str(&format!(" {u} nwarps_ = {nthreads} >> 5;\n"));
s.push_str(&format!(" {ptr} arena = scratch + ((size_t)warp_ * 32u + lane) * {words_per_lane}u;\n"));
s.push_str(&format!(" for ({u} g_ = warp_; g_ < groups; g_ += nwarps_) {{\n"));
s.push_str(&format!(" {u} gid = g_ * 32u + lane;\n"));
s.push_str(&format!(" {u} gbase = baseNonce + g_ * 32u;\n"));
s.push_str(&format!(" {u} tag = salt + g_;\n"));
s
let mut l = String::new();
l.push_str(&format!(" for ({u} g_ = warp_; g_ < groups; g_ += nwarps_) {{\n"));
l.push_str(&format!(" {u} gid = g_ * 32u + lane;\n"));
l.push_str(&format!(" {u} gbase = baseNonce + g_ * 32u;\n"));
l.push_str(&format!(" {u} tag = salt + g_;\n"));
(s, l)
}
/// The scratch lines of program.h (variant 5).
@ -875,21 +884,36 @@ pub fn opencl_kernel_bound(p: &Program, memhard: Option<&MixParams>) -> String {
);
let scratch_args = if p.has_scratch() { ", __global uint* scratch, uint groups, uint salt" } else { "" };
s.push_str(&format!("IGNEUM_KERNEL_HASH void igneum_hash_bound(__global const uint* ds, __global ulong* out, uint baseNonce, uint mask, __global const uint* initw{scratch_args}) {{\n"));
let (setup, unit_loop) = persistent_prologue_parts(CoreDialect::OpenCl, p.class.scratch_words_per_lane());
if p.has_scratch() {
s.push_str(&persistent_prologue(CoreDialect::OpenCl, p.class.scratch_words_per_lane()));
s.push_str(&setup);
} else {
s.push_str(" uint gid = (uint)get_global_id(0);\n");
}
s.push_str(" uint lid = (uint)get_local_id(0);\n");
s.push_str(" uint nonce = baseNonce + gid;\n");
s.push_str(" uint r0, r1, r2, r3, r4, r5, r6, r7;\n");
s.push_str(" uint iw0 = initw[0], iw1 = initw[1], iw2 = initw[2], iw3 = initw[3], iw4 = initw[4], iw5 = initw[5], iw6 = initw[6], iw7 = initw[7];\n");
s.push_str("#if IGNEUM_EXCHANGE == 0\n");
s.push_str(" IGNEUM_LOCAL_WORDS(xch, 2 * IGNEUM_GROUP);\n");
s.push_str(" uint xk = 0u;\n");
s.push_str("#else\n");
s.push_str(" (void)lid;\n");
s.push_str("#endif\n");
if p.has_scratch() {
// the __local exchange buffer must sit at the kernel's outermost scope: declare it, then open the unit loop
s.push_str("#if IGNEUM_EXCHANGE == 0\n");
s.push_str(" IGNEUM_LOCAL_WORDS(xch, 2 * IGNEUM_GROUP);\n");
s.push_str(" uint xk = 0u;\n");
s.push_str("#else\n");
s.push_str(" (void)lid;\n");
s.push_str("#endif\n");
s.push_str(&unit_loop);
s.push_str(" uint nonce = baseNonce + gid;\n");
s.push_str(" uint r0, r1, r2, r3, r4, r5, r6, r7;\n");
s.push_str(" uint iw0 = initw[0], iw1 = initw[1], iw2 = initw[2], iw3 = initw[3], iw4 = initw[4], iw5 = initw[5], iw6 = initw[6], iw7 = initw[7];\n");
} else {
s.push_str(" uint nonce = baseNonce + gid;\n");
s.push_str(" uint r0, r1, r2, r3, r4, r5, r6, r7;\n");
s.push_str(" uint iw0 = initw[0], iw1 = initw[1], iw2 = initw[2], iw3 = initw[3], iw4 = initw[4], iw5 = initw[5], iw6 = initw[6], iw7 = initw[7];\n");
s.push_str("#if IGNEUM_EXCHANGE == 0\n");
s.push_str(" IGNEUM_LOCAL_WORDS(xch, 2 * IGNEUM_GROUP);\n");
s.push_str(" uint xk = 0u;\n");
s.push_str("#else\n");
s.push_str(" (void)lid;\n");
s.push_str("#endif\n");
}
if p.has_wide() {
s.push_str(" uint lane = lid & 31u;\n uint wmask = mask & ~31u;\n");
}
@ -1007,20 +1031,33 @@ pub fn opencl_kernel(p: &Program, memhard: Option<&MixParams>) -> String {
s.push_str(&scratch_prelude(p, CoreDialect::OpenCl));
let scratch_args = if p.has_scratch() { ", __global uint* scratch, uint groups, uint salt" } else { "" };
s.push_str(&format!("IGNEUM_KERNEL_HASH void igneum_hash(__global const uint* ds, __global ulong* out, uint baseNonce, uint mask{scratch_args}) {{\n"));
let (setup, unit_loop) = persistent_prologue_parts(CoreDialect::OpenCl, p.class.scratch_words_per_lane());
if p.has_scratch() {
s.push_str(&persistent_prologue(CoreDialect::OpenCl, p.class.scratch_words_per_lane()));
s.push_str(&setup);
} else {
s.push_str(" uint gid = (uint)get_global_id(0);\n");
}
s.push_str(" uint lid = (uint)get_local_id(0);\n");
s.push_str(" uint nonce = baseNonce + gid;\n");
s.push_str(" uint r0, r1, r2, r3, r4, r5, r6, r7;\n");
s.push_str("#if IGNEUM_EXCHANGE == 0\n");
s.push_str(" IGNEUM_LOCAL_WORDS(xch, 2 * IGNEUM_GROUP);\n");
s.push_str(" uint xk = 0u;\n");
s.push_str("#else\n");
s.push_str(" (void)lid;\n");
s.push_str("#endif\n");
if p.has_scratch() {
s.push_str("#if IGNEUM_EXCHANGE == 0\n");
s.push_str(" IGNEUM_LOCAL_WORDS(xch, 2 * IGNEUM_GROUP);\n");
s.push_str(" uint xk = 0u;\n");
s.push_str("#else\n");
s.push_str(" (void)lid;\n");
s.push_str("#endif\n");
s.push_str(&unit_loop);
s.push_str(" uint nonce = baseNonce + gid;\n");
s.push_str(" uint r0, r1, r2, r3, r4, r5, r6, r7;\n");
} else {
s.push_str(" uint nonce = baseNonce + gid;\n");
s.push_str(" uint r0, r1, r2, r3, r4, r5, r6, r7;\n");
s.push_str("#if IGNEUM_EXCHANGE == 0\n");
s.push_str(" IGNEUM_LOCAL_WORDS(xch, 2 * IGNEUM_GROUP);\n");
s.push_str(" uint xk = 0u;\n");
s.push_str("#else\n");
s.push_str(" (void)lid;\n");
s.push_str("#endif\n");
}
if p.has_wide() {
s.push_str(" uint lane = lid & 31u;\n uint wmask = mask & ~31u;\n");
}

12
proto-cuda/packs-readwidth/scr0k32/kernel.cl
View file

@ -186,19 +186,19 @@ IGNEUM_KERNEL_HASH void igneum_hash(__global const uint* ds, __global ulong* out
uint warp_ = (uint)get_global_id(0) >> 5;
uint nwarps_ = (uint)get_global_size(0) >> 5;
__global uint* arena = scratch + ((size_t)warp_ * 32u + lane) * 256u;
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
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
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
{ 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]

26
proto-cuda/packs-readwidth/scr0k32/kernel_bound.cl
View file

@ -186,19 +186,19 @@ IGNEUM_KERNEL_HASH void igneum_hash(__global const uint* ds, __global ulong* out
uint warp_ = (uint)get_global_id(0) >> 5;
uint nwarps_ = (uint)get_global_size(0) >> 5;
__global uint* arena = scratch + ((size_t)warp_ * 32u + lane) * 256u;
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
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
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
{ 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]
@ -296,20 +296,20 @@ IGNEUM_KERNEL_HASH void igneum_hash_bound(__global const uint* ds, __global ulon
uint warp_ = (uint)get_global_id(0) >> 5;
uint nwarps_ = (uint)get_global_size(0) >> 5;
__global uint* arena = scratch + ((size_t)warp_ * 32u + lane) * 256u;
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint lid = (uint)get_local_id(0);
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
uint iw0 = initw[0], iw1 = initw[1], iw2 = initw[2], iw3 = initw[3], iw4 = initw[4], iw5 = initw[5], iw6 = initw[6], iw7 = initw[7];
#if IGNEUM_EXCHANGE == 0
IGNEUM_LOCAL_WORDS(xch, 2 * IGNEUM_GROUP);
uint xk = 0u;
#else
(void)lid;
#endif
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
uint iw0 = initw[0], iw1 = initw[1], iw2 = initw[2], iw3 = initw[3], iw4 = initw[4], iw5 = initw[5], iw6 = initw[6], iw7 = initw[7];
{ uint x = nonce ^ iw0; x += 0x9e3779b9u * 1u; x = splitmix32(x); r0 = x ^ iw1; }
{ uint x = nonce ^ iw1; x += 0x9e3779b9u * 2u; x = splitmix32(x); r1 = x ^ iw2; }
{ uint x = nonce ^ iw2; x += 0x9e3779b9u * 3u; x = splitmix32(x); r2 = x ^ iw3; }

12
proto-cuda/packs-readwidth/scr2k128/kernel.cl
View file

@ -186,19 +186,19 @@ IGNEUM_KERNEL_HASH void igneum_hash(__global const uint* ds, __global ulong* out
uint warp_ = (uint)get_global_id(0) >> 5;
uint nwarps_ = (uint)get_global_size(0) >> 5;
__global uint* arena = scratch + ((size_t)warp_ * 32u + lane) * 1024u;
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
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
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
{ 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]

26
proto-cuda/packs-readwidth/scr2k128/kernel_bound.cl
View file

@ -186,19 +186,19 @@ IGNEUM_KERNEL_HASH void igneum_hash(__global const uint* ds, __global ulong* out
uint warp_ = (uint)get_global_id(0) >> 5;
uint nwarps_ = (uint)get_global_size(0) >> 5;
__global uint* arena = scratch + ((size_t)warp_ * 32u + lane) * 1024u;
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
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
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
{ 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]
@ -296,20 +296,20 @@ IGNEUM_KERNEL_HASH void igneum_hash_bound(__global const uint* ds, __global ulon
uint warp_ = (uint)get_global_id(0) >> 5;
uint nwarps_ = (uint)get_global_size(0) >> 5;
__global uint* arena = scratch + ((size_t)warp_ * 32u + lane) * 1024u;
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint lid = (uint)get_local_id(0);
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
uint iw0 = initw[0], iw1 = initw[1], iw2 = initw[2], iw3 = initw[3], iw4 = initw[4], iw5 = initw[5], iw6 = initw[6], iw7 = initw[7];
#if IGNEUM_EXCHANGE == 0
IGNEUM_LOCAL_WORDS(xch, 2 * IGNEUM_GROUP);
uint xk = 0u;
#else
(void)lid;
#endif
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
uint iw0 = initw[0], iw1 = initw[1], iw2 = initw[2], iw3 = initw[3], iw4 = initw[4], iw5 = initw[5], iw6 = initw[6], iw7 = initw[7];
{ uint x = nonce ^ iw0; x += 0x9e3779b9u * 1u; x = splitmix32(x); r0 = x ^ iw1; }
{ uint x = nonce ^ iw1; x += 0x9e3779b9u * 2u; x = splitmix32(x); r1 = x ^ iw2; }
{ uint x = nonce ^ iw2; x += 0x9e3779b9u * 3u; x = splitmix32(x); r2 = x ^ iw3; }

12
proto-cuda/packs-readwidth/scr2k32/kernel.cl
View file

@ -186,19 +186,19 @@ IGNEUM_KERNEL_HASH void igneum_hash(__global const uint* ds, __global ulong* out
uint warp_ = (uint)get_global_id(0) >> 5;
uint nwarps_ = (uint)get_global_size(0) >> 5;
__global uint* arena = scratch + ((size_t)warp_ * 32u + lane) * 256u;
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
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
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
{ 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]

26
proto-cuda/packs-readwidth/scr2k32/kernel_bound.cl
View file

@ -186,19 +186,19 @@ IGNEUM_KERNEL_HASH void igneum_hash(__global const uint* ds, __global ulong* out
uint warp_ = (uint)get_global_id(0) >> 5;
uint nwarps_ = (uint)get_global_size(0) >> 5;
__global uint* arena = scratch + ((size_t)warp_ * 32u + lane) * 256u;
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
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
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
{ 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]
@ -296,20 +296,20 @@ IGNEUM_KERNEL_HASH void igneum_hash_bound(__global const uint* ds, __global ulon
uint warp_ = (uint)get_global_id(0) >> 5;
uint nwarps_ = (uint)get_global_size(0) >> 5;
__global uint* arena = scratch + ((size_t)warp_ * 32u + lane) * 256u;
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint lid = (uint)get_local_id(0);
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
uint iw0 = initw[0], iw1 = initw[1], iw2 = initw[2], iw3 = initw[3], iw4 = initw[4], iw5 = initw[5], iw6 = initw[6], iw7 = initw[7];
#if IGNEUM_EXCHANGE == 0
IGNEUM_LOCAL_WORDS(xch, 2 * IGNEUM_GROUP);
uint xk = 0u;
#else
(void)lid;
#endif
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
uint iw0 = initw[0], iw1 = initw[1], iw2 = initw[2], iw3 = initw[3], iw4 = initw[4], iw5 = initw[5], iw6 = initw[6], iw7 = initw[7];
{ uint x = nonce ^ iw0; x += 0x9e3779b9u * 1u; x = splitmix32(x); r0 = x ^ iw1; }
{ uint x = nonce ^ iw1; x += 0x9e3779b9u * 2u; x = splitmix32(x); r1 = x ^ iw2; }
{ uint x = nonce ^ iw2; x += 0x9e3779b9u * 3u; x = splitmix32(x); r2 = x ^ iw3; }

12
proto-cuda/packs-readwidth/scr4k128/kernel.cl
View file

@ -186,19 +186,19 @@ IGNEUM_KERNEL_HASH void igneum_hash(__global const uint* ds, __global ulong* out
uint warp_ = (uint)get_global_id(0) >> 5;
uint nwarps_ = (uint)get_global_size(0) >> 5;
__global uint* arena = scratch + ((size_t)warp_ * 32u + lane) * 1024u;
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
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
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
{ 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]

26
proto-cuda/packs-readwidth/scr4k128/kernel_bound.cl
View file

@ -186,19 +186,19 @@ IGNEUM_KERNEL_HASH void igneum_hash(__global const uint* ds, __global ulong* out
uint warp_ = (uint)get_global_id(0) >> 5;
uint nwarps_ = (uint)get_global_size(0) >> 5;
__global uint* arena = scratch + ((size_t)warp_ * 32u + lane) * 1024u;
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
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
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
{ 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]
@ -296,20 +296,20 @@ IGNEUM_KERNEL_HASH void igneum_hash_bound(__global const uint* ds, __global ulon
uint warp_ = (uint)get_global_id(0) >> 5;
uint nwarps_ = (uint)get_global_size(0) >> 5;
__global uint* arena = scratch + ((size_t)warp_ * 32u + lane) * 1024u;
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint lid = (uint)get_local_id(0);
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
uint iw0 = initw[0], iw1 = initw[1], iw2 = initw[2], iw3 = initw[3], iw4 = initw[4], iw5 = initw[5], iw6 = initw[6], iw7 = initw[7];
#if IGNEUM_EXCHANGE == 0
IGNEUM_LOCAL_WORDS(xch, 2 * IGNEUM_GROUP);
uint xk = 0u;
#else
(void)lid;
#endif
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
uint iw0 = initw[0], iw1 = initw[1], iw2 = initw[2], iw3 = initw[3], iw4 = initw[4], iw5 = initw[5], iw6 = initw[6], iw7 = initw[7];
{ uint x = nonce ^ iw0; x += 0x9e3779b9u * 1u; x = splitmix32(x); r0 = x ^ iw1; }
{ uint x = nonce ^ iw1; x += 0x9e3779b9u * 2u; x = splitmix32(x); r1 = x ^ iw2; }
{ uint x = nonce ^ iw2; x += 0x9e3779b9u * 3u; x = splitmix32(x); r2 = x ^ iw3; }

12
proto-cuda/packs-readwidth/scr4k32/kernel.cl
View file

@ -186,19 +186,19 @@ IGNEUM_KERNEL_HASH void igneum_hash(__global const uint* ds, __global ulong* out
uint warp_ = (uint)get_global_id(0) >> 5;
uint nwarps_ = (uint)get_global_size(0) >> 5;
__global uint* arena = scratch + ((size_t)warp_ * 32u + lane) * 256u;
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
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
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
{ 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]

26
proto-cuda/packs-readwidth/scr4k32/kernel_bound.cl
View file

@ -186,19 +186,19 @@ IGNEUM_KERNEL_HASH void igneum_hash(__global const uint* ds, __global ulong* out
uint warp_ = (uint)get_global_id(0) >> 5;
uint nwarps_ = (uint)get_global_size(0) >> 5;
__global uint* arena = scratch + ((size_t)warp_ * 32u + lane) * 256u;
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
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
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
{ 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]
@ -296,20 +296,20 @@ IGNEUM_KERNEL_HASH void igneum_hash_bound(__global const uint* ds, __global ulon
uint warp_ = (uint)get_global_id(0) >> 5;
uint nwarps_ = (uint)get_global_size(0) >> 5;
__global uint* arena = scratch + ((size_t)warp_ * 32u + lane) * 256u;
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint lid = (uint)get_local_id(0);
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
uint iw0 = initw[0], iw1 = initw[1], iw2 = initw[2], iw3 = initw[3], iw4 = initw[4], iw5 = initw[5], iw6 = initw[6], iw7 = initw[7];
#if IGNEUM_EXCHANGE == 0
IGNEUM_LOCAL_WORDS(xch, 2 * IGNEUM_GROUP);
uint xk = 0u;
#else
(void)lid;
#endif
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
uint iw0 = initw[0], iw1 = initw[1], iw2 = initw[2], iw3 = initw[3], iw4 = initw[4], iw5 = initw[5], iw6 = initw[6], iw7 = initw[7];
{ uint x = nonce ^ iw0; x += 0x9e3779b9u * 1u; x = splitmix32(x); r0 = x ^ iw1; }
{ uint x = nonce ^ iw1; x += 0x9e3779b9u * 2u; x = splitmix32(x); r1 = x ^ iw2; }
{ uint x = nonce ^ iw2; x += 0x9e3779b9u * 3u; x = splitmix32(x); r2 = x ^ iw3; }

12
proto-cuda/packs-readwidth/scr8k128/kernel.cl
View file

@ -186,19 +186,19 @@ IGNEUM_KERNEL_HASH void igneum_hash(__global const uint* ds, __global ulong* out
uint warp_ = (uint)get_global_id(0) >> 5;
uint nwarps_ = (uint)get_global_size(0) >> 5;
__global uint* arena = scratch + ((size_t)warp_ * 32u + lane) * 1024u;
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
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
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
{ 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]

26
proto-cuda/packs-readwidth/scr8k128/kernel_bound.cl
View file

@ -186,19 +186,19 @@ IGNEUM_KERNEL_HASH void igneum_hash(__global const uint* ds, __global ulong* out
uint warp_ = (uint)get_global_id(0) >> 5;
uint nwarps_ = (uint)get_global_size(0) >> 5;
__global uint* arena = scratch + ((size_t)warp_ * 32u + lane) * 1024u;
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
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
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
{ 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]
@ -296,20 +296,20 @@ IGNEUM_KERNEL_HASH void igneum_hash_bound(__global const uint* ds, __global ulon
uint warp_ = (uint)get_global_id(0) >> 5;
uint nwarps_ = (uint)get_global_size(0) >> 5;
__global uint* arena = scratch + ((size_t)warp_ * 32u + lane) * 1024u;
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint lid = (uint)get_local_id(0);
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
uint iw0 = initw[0], iw1 = initw[1], iw2 = initw[2], iw3 = initw[3], iw4 = initw[4], iw5 = initw[5], iw6 = initw[6], iw7 = initw[7];
#if IGNEUM_EXCHANGE == 0
IGNEUM_LOCAL_WORDS(xch, 2 * IGNEUM_GROUP);
uint xk = 0u;
#else
(void)lid;
#endif
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
uint iw0 = initw[0], iw1 = initw[1], iw2 = initw[2], iw3 = initw[3], iw4 = initw[4], iw5 = initw[5], iw6 = initw[6], iw7 = initw[7];
{ uint x = nonce ^ iw0; x += 0x9e3779b9u * 1u; x = splitmix32(x); r0 = x ^ iw1; }
{ uint x = nonce ^ iw1; x += 0x9e3779b9u * 2u; x = splitmix32(x); r1 = x ^ iw2; }
{ uint x = nonce ^ iw2; x += 0x9e3779b9u * 3u; x = splitmix32(x); r2 = x ^ iw3; }

12
proto-cuda/packs-readwidth/scr8k32/kernel.cl
View file

@ -186,19 +186,19 @@ IGNEUM_KERNEL_HASH void igneum_hash(__global const uint* ds, __global ulong* out
uint warp_ = (uint)get_global_id(0) >> 5;
uint nwarps_ = (uint)get_global_size(0) >> 5;
__global uint* arena = scratch + ((size_t)warp_ * 32u + lane) * 256u;
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
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
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
{ 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]

26
proto-cuda/packs-readwidth/scr8k32/kernel_bound.cl
View file

@ -186,19 +186,19 @@ IGNEUM_KERNEL_HASH void igneum_hash(__global const uint* ds, __global ulong* out
uint warp_ = (uint)get_global_id(0) >> 5;
uint nwarps_ = (uint)get_global_size(0) >> 5;
__global uint* arena = scratch + ((size_t)warp_ * 32u + lane) * 256u;
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
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
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
{ 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]
@ -296,20 +296,20 @@ IGNEUM_KERNEL_HASH void igneum_hash_bound(__global const uint* ds, __global ulon
uint warp_ = (uint)get_global_id(0) >> 5;
uint nwarps_ = (uint)get_global_size(0) >> 5;
__global uint* arena = scratch + ((size_t)warp_ * 32u + lane) * 256u;
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint lid = (uint)get_local_id(0);
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
uint iw0 = initw[0], iw1 = initw[1], iw2 = initw[2], iw3 = initw[3], iw4 = initw[4], iw5 = initw[5], iw6 = initw[6], iw7 = initw[7];
#if IGNEUM_EXCHANGE == 0
IGNEUM_LOCAL_WORDS(xch, 2 * IGNEUM_GROUP);
uint xk = 0u;
#else
(void)lid;
#endif
for (uint g_ = warp_; g_ < groups; g_ += nwarps_) {
uint gid = g_ * 32u + lane;
uint gbase = baseNonce + g_ * 32u;
uint tag = salt + g_;
uint nonce = baseNonce + gid;
uint r0, r1, r2, r3, r4, r5, r6, r7;
uint iw0 = initw[0], iw1 = initw[1], iw2 = initw[2], iw3 = initw[3], iw4 = initw[4], iw5 = initw[5], iw6 = initw[6], iw7 = initw[7];
{ uint x = nonce ^ iw0; x += 0x9e3779b9u * 1u; x = splitmix32(x); r0 = x ^ iw1; }
{ uint x = nonce ^ iw1; x += 0x9e3779b9u * 2u; x = splitmix32(x); r1 = x ^ iw2; }
{ uint x = nonce ^ iw2; x += 0x9e3779b9u * 3u; x = splitmix32(x); r2 = x ^ iw3; }