From 357ccb9507e5c25ae4db20b419ffe5b835da61f0 Mon Sep 17 00:00:00 2001 From: igneum-labs <337424239+igneum-labs@users.noreply.github.com> Date: Mon, 5 Oct 2026 20:18:35 +0000 Subject: [PATCH] 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 Co-Authored-By: Claude Fable 5.1 --- igneum-pow/src/emit.rs | 85 +++++++++++++------ proto-cuda/packs-readwidth/scr0k32/kernel.cl | 12 +-- .../packs-readwidth/scr0k32/kernel_bound.cl | 26 +++--- proto-cuda/packs-readwidth/scr2k128/kernel.cl | 12 +-- .../packs-readwidth/scr2k128/kernel_bound.cl | 26 +++--- proto-cuda/packs-readwidth/scr2k32/kernel.cl | 12 +-- .../packs-readwidth/scr2k32/kernel_bound.cl | 26 +++--- proto-cuda/packs-readwidth/scr4k128/kernel.cl | 12 +-- .../packs-readwidth/scr4k128/kernel_bound.cl | 26 +++--- proto-cuda/packs-readwidth/scr4k32/kernel.cl | 12 +-- .../packs-readwidth/scr4k32/kernel_bound.cl | 26 +++--- proto-cuda/packs-readwidth/scr8k128/kernel.cl | 12 +-- .../packs-readwidth/scr8k128/kernel_bound.cl | 26 +++--- proto-cuda/packs-readwidth/scr8k32/kernel.cl | 12 +-- .../packs-readwidth/scr8k32/kernel_bound.cl | 26 +++--- 15 files changed, 194 insertions(+), 157 deletions(-) diff --git a/igneum-pow/src/emit.rs b/igneum-pow/src/emit.rs index 77029f280..5c8965b77 100644 --- a/igneum-pow/src/emit.rs +++ b/igneum-pow/src/emit.rs @@ -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"); } diff --git a/proto-cuda/packs-readwidth/scr0k32/kernel.cl b/proto-cuda/packs-readwidth/scr0k32/kernel.cl index 040d671a2..517a52f3f 100644 --- a/proto-cuda/packs-readwidth/scr0k32/kernel.cl +++ b/proto-cuda/packs-readwidth/scr0k32/kernel.cl @@ -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] diff --git a/proto-cuda/packs-readwidth/scr0k32/kernel_bound.cl b/proto-cuda/packs-readwidth/scr0k32/kernel_bound.cl index 057f228d0..1b033b3cb 100644 --- a/proto-cuda/packs-readwidth/scr0k32/kernel_bound.cl +++ b/proto-cuda/packs-readwidth/scr0k32/kernel_bound.cl @@ -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; } diff --git a/proto-cuda/packs-readwidth/scr2k128/kernel.cl b/proto-cuda/packs-readwidth/scr2k128/kernel.cl index 8fc591da5..706ed7e3e 100644 --- a/proto-cuda/packs-readwidth/scr2k128/kernel.cl +++ b/proto-cuda/packs-readwidth/scr2k128/kernel.cl @@ -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] diff --git a/proto-cuda/packs-readwidth/scr2k128/kernel_bound.cl b/proto-cuda/packs-readwidth/scr2k128/kernel_bound.cl index b08bc77a8..314afb38c 100644 --- a/proto-cuda/packs-readwidth/scr2k128/kernel_bound.cl +++ b/proto-cuda/packs-readwidth/scr2k128/kernel_bound.cl @@ -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; } diff --git a/proto-cuda/packs-readwidth/scr2k32/kernel.cl b/proto-cuda/packs-readwidth/scr2k32/kernel.cl index 6befdde06..59e009d77 100644 --- a/proto-cuda/packs-readwidth/scr2k32/kernel.cl +++ b/proto-cuda/packs-readwidth/scr2k32/kernel.cl @@ -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] diff --git a/proto-cuda/packs-readwidth/scr2k32/kernel_bound.cl b/proto-cuda/packs-readwidth/scr2k32/kernel_bound.cl index 9fb87c712..fa26b1c76 100644 --- a/proto-cuda/packs-readwidth/scr2k32/kernel_bound.cl +++ b/proto-cuda/packs-readwidth/scr2k32/kernel_bound.cl @@ -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; } diff --git a/proto-cuda/packs-readwidth/scr4k128/kernel.cl b/proto-cuda/packs-readwidth/scr4k128/kernel.cl index 3917a927d..babcdd9ed 100644 --- a/proto-cuda/packs-readwidth/scr4k128/kernel.cl +++ b/proto-cuda/packs-readwidth/scr4k128/kernel.cl @@ -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] diff --git a/proto-cuda/packs-readwidth/scr4k128/kernel_bound.cl b/proto-cuda/packs-readwidth/scr4k128/kernel_bound.cl index 669eb1d51..93b2e735f 100644 --- a/proto-cuda/packs-readwidth/scr4k128/kernel_bound.cl +++ b/proto-cuda/packs-readwidth/scr4k128/kernel_bound.cl @@ -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; } diff --git a/proto-cuda/packs-readwidth/scr4k32/kernel.cl b/proto-cuda/packs-readwidth/scr4k32/kernel.cl index 00f785968..b1c967eee 100644 --- a/proto-cuda/packs-readwidth/scr4k32/kernel.cl +++ b/proto-cuda/packs-readwidth/scr4k32/kernel.cl @@ -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] diff --git a/proto-cuda/packs-readwidth/scr4k32/kernel_bound.cl b/proto-cuda/packs-readwidth/scr4k32/kernel_bound.cl index 46034e31b..1090d77b8 100644 --- a/proto-cuda/packs-readwidth/scr4k32/kernel_bound.cl +++ b/proto-cuda/packs-readwidth/scr4k32/kernel_bound.cl @@ -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; } diff --git a/proto-cuda/packs-readwidth/scr8k128/kernel.cl b/proto-cuda/packs-readwidth/scr8k128/kernel.cl index 03840512b..09bfaa3c4 100644 --- a/proto-cuda/packs-readwidth/scr8k128/kernel.cl +++ b/proto-cuda/packs-readwidth/scr8k128/kernel.cl @@ -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] diff --git a/proto-cuda/packs-readwidth/scr8k128/kernel_bound.cl b/proto-cuda/packs-readwidth/scr8k128/kernel_bound.cl index 9cb72a38f..ae27d8663 100644 --- a/proto-cuda/packs-readwidth/scr8k128/kernel_bound.cl +++ b/proto-cuda/packs-readwidth/scr8k128/kernel_bound.cl @@ -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; } diff --git a/proto-cuda/packs-readwidth/scr8k32/kernel.cl b/proto-cuda/packs-readwidth/scr8k32/kernel.cl index 54168b39f..39cddf512 100644 --- a/proto-cuda/packs-readwidth/scr8k32/kernel.cl +++ b/proto-cuda/packs-readwidth/scr8k32/kernel.cl @@ -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] diff --git a/proto-cuda/packs-readwidth/scr8k32/kernel_bound.cl b/proto-cuda/packs-readwidth/scr8k32/kernel_bound.cl index 922739e4a..a7356de52 100644 --- a/proto-cuda/packs-readwidth/scr8k32/kernel_bound.cl +++ b/proto-cuda/packs-readwidth/scr8k32/kernel_bound.cl @@ -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; }