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 1/3] 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; } From 2d04a1c553768934e1d6a1b31a96493c031c78e5 Mon Sep 17 00:00:00 2001 From: igneum-labs <337424239+igneum-labs@users.noreply.github.com> Date: Mon, 5 Oct 2026 20:24:40 +0000 Subject: [PATCH 2/3] read-width: bench-log entry and docs/plans/read-width.md (three cards, probes, widths, the per-load mix, the scratch at 32 and 128 KiB, chip model, recommendation: keep v2; w16 the only width that passes the rules and closes nothing); test multiply made wrapping (dev-profile overflow check) Co-Authored-By: Claude Fable 5.1 --- docs/bench-log.md | 47 +++++++++++++++++ docs/plans/read-width.md | 109 +++++++++++++++++++++++++++++++++++++++ igneum-pow/src/verify.rs | 2 +- 3 files changed, 157 insertions(+), 1 deletion(-) create mode 100644 docs/plans/read-width.md diff --git a/docs/bench-log.md b/docs/bench-log.md index 9c0171fc2..d9bb12a0c 100644 --- a/docs/bench-log.md +++ b/docs/bench-log.md @@ -1524,3 +1524,50 @@ What is measured: one BLS12-381 aggregate signature over 16 summed G1 keys plus | on, split 90 s | v3 | 0 / 2 | none / 3 | 278 / 265 | apart | none | 3 on n0 | 2 (n0 reconnected 6 s after the heal, A's chain at about 58 DAA, inside the table) | Reading (the NEW finding, ledger C4). With the module off GHOSTDAG alone converges on the heavier chain and the losing side's records re-determine (F24 works when the chain moves). With the module on the overlay holds during the split (A, with 30% of the frozen table, locks nothing; B locks 7 and 8) and then fails at the heal in the shipped node: B's certificates for blocks off n0's chain are "kept pending until the chain decides (no lock at this index)", n0's chain never decides because GHOSTDAG keeps its heavier tip and nothing turns the certificate into a fork-choice constraint, and once n0's last lock (index 7, DAA 209) is one window old (DAA 329) the frozen table stops applying on A's chain ("no frozen table (no lock on this chain inside the window)"), A's two keys are 100% of A's own window (B's post-cut blocks are red there) and n0 locks 10, 11, 12 alone; B's certificates for 10 and 11 then log CONFLICTING on n0 (n0 log, 17:27:04 to 17:29:54 BST). A finality fork from a 96-s honest partition, no attacker, table intact at the heal; the 150-s run and the v2 control end the same way. The spec's fork choice ("GHOSTDAG among tips through all certified checkpoints", 3.5) is therefore implemented only for certificates over blocks already on the node's chain. Fix named in the ledger entry: verify an off-chain certificate against the table at its own block and let it constrain fork choice (a certificate-driven reorg), then re-determine. Raw: `scratchpad fud-a/c4-results-*.md`, node logs `c4-on90-tmp/`, `c4-v2-control-tmp/`. + +## 5 October 2026 (night), read width of the lottery hash: 4, 16 and 64-byte loads, a per-load mix, a written scratch; three cards (gate 1 experiment, cryptographer) + +Branch `readwidth` (commits 019b014, b970dda, 4badcee, a9e002c, d0018cf and the entry commit); plan and recommendation in `docs/plans/read-width.md`. Nothing here changes consensus: every class sits behind `igneum-pow --class` and the default class is generator version 2 byte for byte (`igneum-pow/tests/packs.rs` passes on the four pinned packs after every commit). Question (the project lead, after "the 9070 XT on the eGPU" above): would wider reads keep the latency-bound random-access property while closing the vendor gap. Additions from the coordinator: a per-load width drawn from an era-fixed mix, and a written per-warp scratch (measurement only, no soundness claim). + +**What a class does** (`igneum-pow/src/generator.rs` `LoadClass`, `verify::fold_words`, the three emitters): a load of W words reads the W-word-aligned address `(src AND MASK) AND NOT (W - 1)` and folds every word into `dst` (`x = dst ^ w0; x = (rotl(x, 11) * 0x9e3779b1) ^ w[j]`); W = 1 is the lottery hash exactly (`w4` = pack `bcc1248b10cc90f2`). A mix class draws W per load with one extra `below(100)` roll per instruction. A scratch class `scrk` turns `k` of the 16 memory slots into read-modify-writes of a 16-byte slot of the lane's share of a `kb` KiB per-warp scratch (kernels run persistent warps, one per block or work-group; a slot reads as a seed-and-base fill until the unit writes it, behind a per-unit tag). Program ids carry the class. Dependent chain and 32-lane unit unchanged. + +**Correctness**: 23 packs (`proto-cuda/packs-readwidth/`, Rust CPU reference vectors). Every pack passed its three vector units and the cache and dataset checks on Metal (M5 Max, `proto-metal/packbench`), Apple OpenCL (`--bench-pack`), the RTX 5090 (NVRTC, `igneum-worker-cuda --bench`) and, the 16 width and mix packs, the RX 9070 XT (`igneum-worker-opencl --bench-pack`); the 2^24 batch fingerprints agree across all four runtimes on every pack (for example w16 `e7c890445b47af60`, w64 `836e56e7d496e980`, mixB-2 `a18ac73098c76007`). The clang CUDA emulation (w16, w64, w64x4, mixA-0, mixB-0: 3 of 3 units standalone and 2 of 2 in batch at 2 warps per block) and the clang OpenCL emulation (the same five plus scr2k32 and scr8k128, sub-group 32 and, width packs, wave64 with sub-group shuffles) pass with equal fingerprints per configuration. Acceptance rule on the classes: 60 candidates per class, rejection 0 to 14 of 60 (w16 and w64 as v2; the mixes the same; the scratch classes' distinct-address bound now covers dataset loads only, since a 64-slot lane scratch repeats slots by design). CPU verifier (M5 Max, one core, avg of 50 units, `igneum-pow bench --class`): v2 0.604 ms, w16 0.610, w64 0.630, w64x4 0.160, mix50-35-15 0.620, mix25-50-25 0.614, scr0k32 0.600 (1.004 on a loaded re-run), scr2k32 0.657, scr4k32 0.458, scr8k32 0.317, scr2k128 0.535, scr4k128 0.458, scr8k128 0.311; per hash divide by 32. The wide reads cost the verifier nothing (a lane's words lie in one item); scratch ops replace item derivations and make it cheaper. + +**Probes** (`--memprobe`, dependent random reads at 1024 MiB, G reads/s, best over lanes in flight; 4 B = the hash's pattern; the 5090 and 9070 XT with the card off in the app, the Mac through Apple OpenCL under a load average of 5 to 10): + +| Card | 4 B chase | 16 B | 64 B | 64 B as GB/s | coalesced stream GB/s (rated) | integer chain | +|---|---|---|---|---|---|---| +| RTX 5090 (PC 2, CUDA) | 17.5 to 18.2 | 18.0 to 19.9 | 9.1 to 15.7 (9.1 at 4 M lanes) | 584 | 1,579 (1,792) | 39.0 T op/s | +| RX 9070 XT (PC 1, eGPU, OpenCL) | 2.42 to 2.66 | 2.43 to 2.73 | 2.47 to 2.87 | 158 | 636 (640) | 6.2 T op/s | +| Apple M5 Max (Apple OpenCL, approximate) | 3.50 | 3.51 | 3.51 | 225 | 522 | | + +Reading: on the 9070 XT and the M5 Max a 64-byte dependent read costs exactly what a 4-byte one costs (the line is fetched either way); on the 5090 a 64-byte read costs about two 4-byte reads (two 32-byte sectors) and the 64 B chase at full occupancy sits at 584 GB/s, a third of the stream. + +**Hash rates** (5 timed dispatches of 2^24 nonces after a warm-up; Metal and the 9070 XT by device time, the 5090 by wall time around the stream sync; the PC cards switched off in the app for the run and restored, PC 1's 5090 and the integrated chip kept mining; the Mac under other agents' builds, load 4 to 9, so its absolute numbers carry that; the share = measured / (the card's probe ceiling at the class's widths / loads per hash)): + +| Class | dataset B/hash | RTX 5090 MH/s (share) | RX 9070 XT MH/s (share) | M5 Max Metal MH/s (share) | 5090 / 9070 | +|---|---|---|---|---|---| +| v2 (w4, the lottery hash) | 512 | 136.1 (0.96) | 18.15 (0.87) | 27.74 (1.01) | 7.5x | +| w16 | 2,048 | 139.8 (0.90) | 17.90 (0.84) | 28.26 (1.03) | 7.8x | +| w64 | 8,192 | 71.9 (0.58) | 17.59 (0.78) | 28.27 (1.03) | 4.1x | +| w64x4 (32 loads) | 2,048 | 275.3 (0.56) | 75.19 (0.84) | 109.7 (1.00) | 3.7x | +| mix50-35-15, 6 programs: min / median / max (spread of median) | 1,664 to 3,680 | 99.5 / 114.2 / 121.0 (18.8%) | 17.45 / 18.76 / 18.83 (7.4%) | 25.36 / 27.26 / 28.43 (11.3%) | 6.1x | +| mix25-50-25, 6 programs | 2,240 to 5,024 | 95.9 / 107.3 / 119.8 (22.3%) | 17.84 / 18.45 / 18.85 (5.5%) | 23.21 / 24.68 / 25.21 (8.1%) | 5.8x | + +Scratch (variant 5; N persistent warps; 5090: 2,048 warps launched against a resident capacity of 4,080 = 24 blocks/SM x 1 warp/block x 170 SMs at `--block-warps 1`, the occupancy query unchanged by the allocation (24 before and after); Metal: 2,048 to 16,384 warps swept, best shown; arena = N x per-warp size; the whole working set = 1 GiB dataset + 256 MiB cache + 128 MiB output + arena, under 2 GB on every row): + +| Class (k of 16 slots, KiB per warp) | scratch ops/hash | dataset B/hash | RTX 5090 MH/s (vs scr0, share) | M5 Max Metal MH/s (vs scr0) | RX 9070 XT MH/s | 5090 arena / working set | +|---|---|---|---|---|---|---| +| scr0k32 (control, persistent loop, no RMW) | 0 | 512 | 139.1 (0, 0.98) | 28.25 (0) | 17.88 (control, 0.86) | 64 MiB / 1.4 GiB | +| scr2k32 (12.5%) | 16 | 448 | 114.4 (-18%, 0.80) | 26.14 (-7%) | 14.65 (-18%) | 64 MiB / 1.4 GiB | +| scr4k32 (25%) | 32 | 384 | 109.8 (-21%, 0.76) | 31.74 (+12%) | 14.00 (-22%) | 64 MiB / 1.4 GiB | +| scr8k32 (50%) | 64 | 256 | 122.1 (-12%, 0.82) | 49.08 (+74%) | 14.17 (-21%) | 64 MiB / 1.4 GiB | +| scr2k128 (12.5%) | 16 | 448 | 110.1 (-21%, 0.77) | 26.24 (-7%) | 14.07 (-21%) | 256 MiB / 1.6 GiB | +| scr4k128 (25%) | 32 | 384 | 98.0 (-30%, 0.68) | 28.08 (-1%) | 13.14 (-27%) | 256 MiB / 1.6 GiB | +| scr8k128 (50%) | 64 | 256 | 72.8 (-48%, 0.49) | 35.44 (+25%) | 12.03 (-33%) | 256 MiB / 1.6 GiB | + +The 9070 XT rows are 2,048 persistent warps (4,096 within 1 percent), arena 64 MiB at 32 KiB and 256 MiB at 128 KiB, working set 1.4 and 1.6 GiB; its control (17.88, the persistent loop) equals its v2 rate (18.15) within 2 percent, and every RMW share costs it 18 to 33 percent: on AMD a scratch op is a dependent 16-byte read plus a write into a region the 64 MB Infinity Cache does not hold for 2,048 warps, so it is memory work there as on the 5090, not the cached op it is on Apple. Apple OpenCL on the same scratch packs (wall time, `--bench-pack --warps 2048`): scr0k32 27.85, scr2k32 28.58, scr4k32 32.43, scr8k32 47.93, scr2k128 25.67, scr4k128 27.24, scr8k128 32.93 MH/s, the Metal shape within 4 percent, fingerprints equal. Bytes moved per scratch op: 16 read + 16 written (the tag word included); per hash at 50 percent, 1,024 read + 1,024 written beside 256 of dataset reads. The 5090 at 4,096 launched warps (above its 4,080 resident) lost 2 to 26 percent (scr8k32 90.0 MH/s), so the rows above are the in-capacity launch. + +**Readings.** (1) Same count, wider: the vendor gap does not move at 16 B (7.8x) because on the 9070 XT a 4-byte read already costs a 64-byte line and on the 5090 a 16-byte read costs one 32-byte sector, the same as 4 bytes: the memory systems do identical work, only the fold's input grows. At 64 B the gap closes to 4.1x, entirely by the 5090 losing half its rate (its share falls to 0.58 and its DRAM traffic reaches 589 GB/s, 37 percent of the stream: bandwidth, not latency, bounds it), while the 9070 XT and the M5 Max do not move. (2) Fewer, wider (w64x4): 3.7x, but every card runs 4x faster because the dependent chain is 32 loads long instead of 128; the 5090 sits at a 0.56 share (bandwidth), so a chip with more bandwidth per dollar than a GPU gains, which is the Ethash shape the design avoids. (3) The mix: the hour-to-hour spread is 7 to 22 percent of the median per card (the 5090 the widest, because its 64-byte loads are the expensive ones and their count per program runs 2 to 8 of 16); the programs with many 64-byte loads (mixA-3, mixA-5, mixB-2) are the slow hours on the 5090 and the fast ones nowhere. (4) The scratch: on the 5090 every RMW share costs 12 to 48 percent against the persistent control, the 32 KiB arena less than the 128 KiB one (the smaller arena, 64 MiB over 2,048 warps, sits inside the 96 MB L2); on the M5 Max the 32 KiB rows are FASTER than the control (+12 and +74 percent at 25 and 50 percent), because the arena (128 MiB over 4,096 warps) lives in the chip's caches and a scratch op is cheaper than a dataset read, so replacing dataset loads raises the rate: the scratch at these sizes is not memory work on Apple and is partly cached on NVIDIA. The chip row for these variants comes from the ca2-soundness branch; what this entry gives is the GPU cost and the share. (5) Latency-bound shares: v2 0.87 to 1.01 on the three cards, w16 0.84 to 1.03, w64 0.58 (5090) and 0.78 (9070 XT); the Mac's shares above 1 are an Apple OpenCL probe under load against a Metal rate. + +Jobs: `run-readwidth-5090-20261005` and `run-readwidth-9070-20261005` (probes; the packs refused for their string seeds, fixed in a9e002c), `run-readwidth-5090-20261005c`, `run-readwidth-9070-20261005c` (benches), `run-readwidth-9070-scratch-20261005d` (the scratch packs after the `__local` fix d0018cf, AMD's compiler requires the exchange buffer at the kernel's outermost scope); read back with `node tools/jobs.mjs --all`. Mac commands and logs: `docs/plans/read-width.md` section 3. The worker exes for the jobs: `proto-cuda/nvrtc/build-windows.sh` on this branch (mingw), sha256 of the CUDA one `6f46336f...defe1`. diff --git a/docs/plans/read-width.md b/docs/plans/read-width.md new file mode 100644 index 000000000..60a4d3da8 --- /dev/null +++ b/docs/plans/read-width.md @@ -0,0 +1,109 @@ +# Read width of the lottery hash: 4, 16 and 64-byte loads, a per-load mix, and a written scratch (gate 1 experiment) + +5 October 2026. Branch `readwidth` (worktree `../igneum-wt-readwidth`), commits 019b014 and b970dda plus the measurement commit. Nothing here changes consensus, the live generator, the pinned vectors or a shipped binary: every class sits behind `--class` in `igneum-pow` and the default class is generator version 2 byte for byte (`igneum-pow/tests/packs.rs` still compares the four pinned packs against the emitters). Numbers and a recommendation; the decision is the project lead's. + +## 1. The question + +The bench-log entry "the 9070 XT on the eGPU" (5 October 2026) found the hash bound by dependent random 4-byte reads over the 1 GiB dataset, 128 per hash: the RX 9070 XT finishes 2.4 to 2.7 G such reads a second (18 MH/s), the RTX 5090 16 to 18 G (127 MH/s), the M5 Max 3.45 G (23 to 28 MH/s). AMD fetches a 64-byte line per 4-byte read, so 94 percent of its memory traffic is unused; NVIDIA fetches a 32-byte sector and its 96 MB L2 catches a share. the project lead's question: would wider reads keep the chip-resistance property (latency-bound, random access) while closing the vendor gap? Two additions from the coordinator: a per-load width drawn from an era-fixed mix so no chip is built for one width, and a written per-warp scratch so part of the memory work cannot be mirrored into read-only SRAM. + +## 2. What was built (all behind the flag) + +| Class (`--class`) | Loads per hash | What a load does | Dataset bytes per hash | +|---|---|---|---| +| `v2` (= `w4`, the lottery hash) | 128 | `dst ^= dataset[src & MASK]`, one 4-byte word | 512 | +| `w16` | 128 | the 16-byte-aligned group of 4 words at `src & MASK`, every word folded into `dst` | 2,048 | +| `w64` | 128 | the 64-byte-aligned item (16 words), every word folded | 8,192 | +| `w64x4` | 32 (4 load slots) | as `w64`; the same bytes per hash as 512 loads of 4 bytes | 2,048 | +| `mix50-35-15` | 128 | per load, width 4, 16 or 64 bytes drawn from the program stream with probabilities 50/35/15 | 1,664 to 3,680 over the six programs measured (expected 2,202) | +| `mix25-50-25` | 128 | the same with 25/50/25 | 2,240 to 5,024 (expected 3,200) | +| `scrk` | 128 memory operations | `k` of the 16 slots are scratch read-modify-writes into a `kb` KiB per-warp scratch (16-byte slots, lane-major); the other `16 - k` are 4-byte loads | 4 x (16 - k) x 8 reads plus 16 B read and 16 B written per scratch op | + +The fold. A load of W words reads the W-word-aligned address `b = (src AND MASK) AND NOT (W - 1)` and sets `x = dst XOR w[0]; for j in 1..W: x = (rotl(x, 11) * 0x9e3779b1) XOR w[j]; dst = x` (`verify::fold_words`, mirrored in the three kernel dialects). The rotate-multiply between the words makes the fold state-dependent: two different lines give two different maps of `dst` (the multiply by an odd constant is not xor-linear), so no function of the line alone can stand in for it, and a dataset of pre-folded lines cannot replace the dataset. The dependent chain is unchanged: the next load's address comes from a register that the fold wrote. Width 1 is the lottery hash's xor of one word, so `w4` is the pinned program `bcc1248b10cc90f2` exactly. + +The per-load draw. Every class other than `v2` takes one extra draw per instruction (`below(100)`, the width roll, consumed on every slot so the stream stays uniform) after the nine draws of spec 01 section 1.4.3; on a load slot the width is the first entry of the mix whose cumulative weight exceeds the roll. The program id of a class is `FNV-1a-64("igneum-program-rw/" || 2 || seed words || attempt || mix[3] || slots [|| "scratch/" k kb])`, so no class program can pass for a version 2 program. The acceptance rule of 1.4.6 runs unchanged on the aligned addresses (lane-constant sites and the distinct-address bound, scaled to the dataset loads per hash). + +The scratch (variant 5, measurement only). The kernel runs N persistent warps (one per block or work-group of 32); warp `w` owns scratch `w` and runs units `w, w + N, ...` of the launch. A scratch op reads the lane's 16-byte slot `src AND (slots - 1)`: three data words behind a per-unit tag; a slot whose tag is not this unit's reads as its fill `splitmix32(((base + lane) XOR seed[j]) + slot x 0x9e3779b1 + (j + 1) x 0x85ebca77)`, the three words are folded into `dst` as above, and the slot is rewritten `(tag, x XOR w1, rotl(x, 7) XOR w2, x + w0)`. The CPU verifier holds the touched slots of one unit (at most 32 x k x 8) and nothing else. The scratch is per lane (a 32 KiB warp scratch is 64 slots per lane, 128 KiB is 256), so two lanes never race on a slot and the result is a function of (program, day, unit) alone; the GPU's tags make the lazy fill exact as long as a tag is not reused within the arena's history (the salt advances per unit; it wraps after 2^32 units, a measurement caveat, not a design). + +## 3. Method + +| Step | Command (every figure in the bench log carries its command) | +|---|---| +| Packs | `igneum-pow export --seed --class --out proto-cuda/packs-readwidth/` (23 packs; the six mix seeds per mix are `igneum-readwidth/A/` and `/B/` with attempt 0 accepted, `A/4` skipped: rejected at attempt 0) | +| CPU verifier | `igneum-pow bench --seed igneum-genesis --class --warps 50` (M5 Max, one core; load average 4 to 9 from other agents' builds during the run) | +| Metal | `proto-metal/packbench --pack --batches 5 --batch-log2 24 --group 256 [--warps N]` (new harness: runs the pack's own text; vectors, cache FNV, dataset words, 2^24 fingerprint, MH/s by GPU time) under the measure lock | +| Apple OpenCL | `proto-opencl/igneum-bench-cl-rw --bench-pack --pack --batches 5 --batch-log2 24` and `--memprobe` (Apple's OpenCL, a correctness check and an approximate rate) | +| Emulators | `proto-cuda/emu/emu.sh ../packs-readwidth/

--batch-log2 13 --batches 1 --block-warps 2` (the CUDA text as C++); `proto-opencl/emu/emu.sh ../packs-readwidth/

0 32 --sg 32` and `1 64 --sg 64` (the OpenCL text, wave32 and wave64) | +| RTX 5090 | job `run-readwidth-5090-20261005` on PC 2 (`relay/playbooks/readwidth-5090.ps1`): the NVIDIA card switched off in the app through `POST app.url/api/cards` and restored after; `igneum-worker-cuda --memprobe`, then `--bench --pack

--batches 5 --batch-log2 24` per pack (NVRTC, the pack's own text, vectors through the bound kernel, 2^24 fingerprint) | +| RX 9070 XT | job `run-readwidth-9070-20261005` on PC 1 (`relay/playbooks/readwidth-9070.ps1`): only the gfx1201 card switched off; `igneum-worker-opencl --device D --memprobe`, then `--bench-pack --pack --batches 5 --batch-log2 24` per pack | + +Latency-bound share = measured MH/s x loads per hash / the card's dependent-read ceiling for that width from its own probe at 1024 MiB (for a mix, the harmonic combination of the widths' ceilings weighted by the program's width counts). A share near 1 means the hash runs at the card's random-access limit, the property the design wants; a share well under 1 means something else bounds it (bandwidth, ALU, occupancy). + +## 4. Results (full tables with commands in `docs/bench-log.md`, "read width of the lottery hash") + +Probe ceilings at 1024 MiB (G dependent reads/s): RTX 5090 4 B 17.5, 16 B 18.0, 64 B 9.1 (584 GB/s), stream 1,579 GB/s; RX 9070 XT 4 B 2.42, 16 B 2.43, 64 B 2.47 (158 GB/s), stream 636; M5 Max (Apple OpenCL, approximate) 3.50 / 3.51 / 3.51, stream 522. + +| Class | dataset B/hash | RTX 5090 MH/s (latency-bound share) | RX 9070 XT (share) | M5 Max Metal (share) | 5090 / 9070 | DRAM bytes moved per hash, NVIDIA 32 B sector / AMD 64 B line | CPU verify ms per unit | +|---|---|---|---|---|---|---|---| +| v2 = w4 (today) | 512 | 136.1 (0.96) | 18.15 (0.87) | 27.74 (1.01) | 7.5x | 4,096 / 8,192 | 0.604 | +| w16 | 2,048 | 139.8 (0.90) | 17.90 (0.84) | 28.26 (1.03) | 7.8x | 4,096 / 8,192 | 0.610 | +| w64 | 8,192 | 71.9 (0.58) | 17.59 (0.78) | 28.27 (1.03) | 4.1x | 8,192 / 8,192 | 0.630 | +| w64x4 (32 loads) | 2,048 | 275.3 (0.56) | 75.19 (0.84) | 109.7 (1.00) | 3.7x | 2,048 / 2,048 | 0.160 | +| mix50-35-15 (6 programs, min / median / max) | 1,664 to 3,680 | 99.5 / 114.2 / 121.0, spread 18.8% | 17.45 / 18.76 / 18.83, 7.4% | 25.36 / 27.26 / 28.43, 11.3% | 6.1x | 5,939 / 8,192 expected | 0.620 | +| mix25-50-25 (6 programs) | 2,240 to 5,024 | 95.9 / 107.3 / 119.8, 22.3% | 17.84 / 18.45 / 18.85, 5.5% | 23.21 / 24.68 / 25.21, 8.1% | 5.8x | 7,168 / 8,192 expected | 0.614 | + +Scratch, variant 5 (N persistent warps; GPU cost against the persistent control scr0k32; working set = 1 GiB + 256 MiB + 128 MiB output + N x size): + +| Class | RMW share | dataset B/hash | scratch B/hash read + written | RTX 5090 MH/s, 2,048 warps of 4,080 resident (vs control, share) | M5 Max Metal (vs control) | RX 9070 XT, 4,096 warps | working set 5090 / 9070 / Mac | +|---|---|---|---|---|---|---|---| +| scr0k32 | 0 | 512 | 0 | 139.1 (control, 0.98) | 28.25 (control) | 17.88 (control, 0.86) | 1.4 GiB / 1.5 GiB / 1.5 GiB | +| scr2k32 | 12.5% | 448 | 256 + 256 | 114.4 (-18%, 0.80) | 26.14 (-7%) | 14.65 (-18%) | same | +| scr4k32 | 25% | 384 | 512 + 512 | 109.8 (-21%, 0.76) | 31.74 (+12%) | 14.00 (-22%) | same | +| scr8k32 | 50% | 256 | 1,024 + 1,024 | 122.1 (-12%, 0.82) | 49.08 (+74%) | 14.17 (-21%) | same | +| scr2k128 | 12.5% | 448 | 256 + 256 | 110.1 (-21%, 0.77) | 26.24 (-7%) | 14.07 (-21%) | 1.6 GiB / 1.9 GiB / 1.9 GiB | +| scr4k128 | 25% | 384 | 512 + 512 | 98.0 (-30%, 0.68) | 28.08 (-1%) | 13.14 (-27%) | same | +| scr8k128 | 50% | 256 | 1,024 + 1,024 | 72.8 (-48%, 0.49) | 35.44 (+25%) | 12.03 (-33%) | same | + +Resident warps and the cap: the 5090 holds 4,080 warps at one warp per block (24 blocks per SM x 170 SMs; 8,160 at 8 warps per block), so 128 KiB each is 510 MiB and the whole working set 1.9 GiB; a 1 MB scratch would have been 4.0 GiB at this geometry and 10.6 GiB at the 64-warp-per-SM figure, which is why the cap moved the size to the tens of kilobytes. The occupancy query returned 24 blocks per SM before and after the arena allocation: the allocation did not change it. The 9070 XT's OpenCL runtime has no occupancy query; 4,096 persistent warps were launched (64 per compute unit over 64 CUs, approximate) and the arena is 128 MiB at 32 KiB, 512 MiB at 128 KiB. The Mac's residency is not reported; 2,048 to 16,384 warps were swept and the best row kept. + +Chip model, re-run with the measured widths (the M16 arithmetic of `docs/analysis/m16-recompute-attacker-2026-10-05.md`; the on-die-cache recompute chip's row per scratch variant is the ca2-soundness branch's, as agreed with the Counter ASIC 2.0 coordinator): + +| Class | what a chip with its own DRAM controller gains over the GPU's memory system | what a chip with on-die SRAM gains | +|---|---|---| +| v2 | the GPU fetches 8 to 16x the bytes it uses (AMD 64 B, NVIDIA 32 B per 4 B); a chip fetching 32 B bursts moves 4,096 B per hash, the 5090's figure, so nothing over NVIDIA and 2x over AMD in traffic, none in latency (the chain is 128 dependent DRAM latencies on either) | the recompute attacker of M16: 150,000 integer ops per hash against the 256 MiB cache; 2.4x at equal silicon before a fixed-function factor (unchanged by the width) | +| w16 | the same: 4,096 / 8,192 bytes moved, 2,048 used; traffic efficiency 50 percent on NVIDIA, 25 on AMD | unchanged: the fold uses every byte, so the chip recomputes 128 items per hash exactly as before; the SRAM mirror of the read-only dataset (1 GiB) stays out of reach | +| w64 | every byte moved is used on both vendors (8,192 moved, 8,192 used); the 5090 is bandwidth-bound at 589 GB/s, so a chip with HBM3 class bandwidth (several TB/s, approximate) is bandwidth-advantaged: the Ethash shape | unchanged in op count; but the chain of 128 loads now moves 8 KB, so a chip's advantage shifts from latency to bandwidth per dollar, which is the wrong direction for the design's 2x target | +| w64x4 | 2,048 moved and used; 32 latencies per hash; every card 4x faster; a bandwidth-rich chip gains as above | 32 items per hash: the recompute attacker's op count falls 4x (37,500 per hash), so the M16 gain rises 4x: fails the 2x target by arithmetic | +| mixes | between v2 and w64 per program; the chip cannot be built for one width, but the GPU pays the 64-byte hours (the 5090 loses up to 27 percent in a heavy hour) | as v2 per item; the recompute attacker is indifferent to the width | +| scratch | a chip must provide writable memory for N units in flight: 32 KiB x N at the GPU's geometry (128 MiB at 4,080), against the 256 MiB read-only cache it could mirror into SRAM (54 to 83 mm^2 at a leading node, the coordinator's figure, approximate); but a unit touches at most k x 8 x 32 slots (4 KiB at 50 percent), the fill is a function and the tags are per unit, so a chip need only hold the touched set per unit in flight (the soundness caveat below) | the dataset reads replaced by scratch ops are reads the chip no longer has to serve from the 1 GiB; at 50 percent the recompute attacker computes 64 items instead of 128 | + +Soundness (variant 5, measurement only, as instructed; the chip row is the ca2-soundness branch's, a465881: the on-die-cache recompute chip's gain is 2.4x at 0, 12.5, 25 and 50 percent, replaced or added, 32 or 128 KB, so the scratch does not move it): the per-unit scratch starts from a fill that any implementation can compute, and a unit writes at most `k x 8` slots per lane; an implementation that keeps only the touched slots of each unit in flight (the CPU verifier does exactly this) needs 16 B x touched slots, not the nominal arena, so the "real memory a chip must provide" is bounded by units in flight x touched slots, not by N x 32 KiB. The variant forces memory that is written, which SRAM can hold as well as DRAM; it does not force memory that is large. A written region that outlives the unit (state carried across units) would, and the CPU verifier could not replay it. This is the finding, not a recommendation. + +## 5. Recommendation (the decision is the project lead's) + +the project lead's rules, as passed by the coordinator: width = the widest read that keeps every card latency-bound (achieved within 90 percent of the probe ceiling at that width) with margin on the 5090 (bytes per hash x rate under a third of the 1,579 GB/s stream); the mix is in only if the six-program spread is under 5 percent per card; the scratch share is the smallest at which the chip model's gain falls under 1.5x at the lowest GPU cost within the 6 GB cap. + +| Variant | Verdict under the rules | Numbers | +|---|---|---| +| w16 (16-byte loads, 128 per hash) | PASSES the rules: shares 0.90 / 0.84 / 1.03 (the 9070 XT's 0.84 equals its v2 share of 0.87 within noise: the card is at its ceiling in both), 286 GB/s on the 5090 = 18 percent of the stream. It does NOT close the vendor gap (7.8x against 7.5x), because the memory systems already move a sector or a line per load; it changes what the fold consumes, nothing the DRAM does | the only width row that passes; a no-cost change in rate (+2.7 percent 5090, -1.4 percent 9070 XT, +1.9 percent M5 Max) | +| w64 | FAILS: 5090 share 0.58, 37 percent of the stream; closes the gap to 4.1x only by making the 5090 bandwidth-bound | | +| w64x4 | FAILS: shares 0.56 / 0.84 / 1.00, the recompute gain rises 4x | | +| mix 50/35/15 and 25/50/25 | OUT: spreads 18.8 and 22.3 percent on the 5090, 7.4 and 5.5 on the 9070 XT, 11.3 and 8.1 on the M5 Max, all over 5 percent; a chip is not built for a width anyway (see the model: the width does not change the recompute attacker) | | +| scratch | OUT: every share costs the 5090 12 to 48 percent and the 9070 XT 18 to 33 percent, and raises the M5 Max's rate (the arena is cached there); the soundness branch's chip row (ca2-soundness a465881, the on-die-cache recompute chip) stays at 2.4x at every share, 32 or 128 KB, because the verifier resets the scratch per unit and the live state is the hash's own read-modify-writes, which a chip keeps in 80 to 320 B per lane; under the rule the share is 0. The rows stay as the measurement that decided it | | + +Recommendation: keep 128 loads per hash and 4 bytes per load (v2) for the devnet; if a width change is wanted for the fold's sake (every byte of the sector consumed, which removes the "94 percent waste" statement from the AMD entry without changing what the card does), w16 is the one that passes every rule and costs nothing measurable, and it is the only width worth a vector re-cut. The AMD gap is a random-access gap (2.4 G against 17.5 G dependent reads per second at 1 GiB on the cards we own), and no read width closes it without turning the 5090 bandwidth-bound; the levers that act on the gap are the ones outside this experiment (the AMD card's memory path, and the dataset size against the 5090's 96 MB L2 share, which the probe's 64 MiB rows show at 9 G reads/s against 2.4 at 1 GiB). The per-load mix is out on stability; the scratch is out on GPU cost and on the soundness caveat. + +## 6. What w16 would change if adopted (not done; the decision is the project lead's) + +| Where | Change | +|---|---| +| `docs/spec/01-lottery-hash.md` 1.4.1 | `load`: `dst = fold(dst, dataset[b .. b + 4))`, `b = (src AND MASK) AND NOT 3`, with the fold written out; 1.4.3 unchanged (no width draw for a fixed width); 1.4.6 unchanged (the aligned address is the address the rule sees) | +| 1.5 | "A load reads one 4-byte word" becomes 16 bytes aligned; the single-form text-search rule of 1.14 item 2 becomes the wide form; the item size (64 B) and `dataset[w] = item(w >> 4)[w AND 15]` unchanged | +| 1.11 | unchanged in count (4,096 items per unit; the verifier derives the same items) | +| 1.15, 1.17 | every vector re-cut (new program ids: the class enters the id or the generator version steps to 3); the four pinned packs replaced; the conformance fuzz re-run on Metal, CUDA and OpenCL (this branch's 23 packs and the three emulators are the template) | +| Litepaper, Mining ("random reads over a multi-gigabyte dataset") and the vs-RandomX "128 dataset addresses" rows | "128 reads of 16 bytes"; `site/bench.html` sector arithmetic (32 B per 4 B) becomes 32 B per 16 B | +| Workers | no host change: the kernel text carries the loads; `proto-cuda/host.cu`'s static mask check (`TESTS.md` section 5) learns the wide form | +| Cost on the 5090 | none measured (+2.7 percent); on the 9070 XT -1.4 percent; verifier +1 percent | + +## 7. Files + +`igneum-pow/src/{generator,verify,accept,emit,memhard,main}.rs` (the classes, behind `--class`), `proto-cuda/packs-readwidth/` (23 packs), `proto-metal/packbench.swift` (Metal from a pack's files), `proto-opencl/host.c` (`--bench-pack`, `--warps`, the 16-byte probe row, the scratch arguments), `proto-cuda/nvrtc/worker.cpp` (`--bench`, `--memprobe`, the scratch arena), `proto-cuda/nvrtc/packfile.h` (class fields; string-seed packs), `proto-cuda/emu/cuda_runtime.h` and `proto-opencl/emu/{emu_opencl.h,emu_main.cpp}` (vector types, the persistent launch), `relay/playbooks/readwidth-*.ps1` (the PC jobs: the card under test off in the app and restored, never the other card). diff --git a/igneum-pow/src/verify.rs b/igneum-pow/src/verify.rs index fedfe1b91..f20bedc60 100644 --- a/igneum-pow/src/verify.rs +++ b/igneum-pow/src/verify.rs @@ -505,7 +505,7 @@ mod tests { let ds = DatasetSource::new("2026-10-03", DatasetMode::MemoryHard, 20); let mut base = [0u32; LANES]; for (k, b) in base.iter_mut().enumerate() { - *b = ((k as u32 * 0x9E37_79B1) & ds.mask) & !15; + *b = ((k as u32).wrapping_mul(0x9E37_79B1) & ds.mask) & !15; } let mut out = [[0u32; 16]; LANES]; let items = ds.fetch_wide(&base, 16, &mut out); From 06dcb310fb5b190e726dbf1e584c94a60379cc4e Mon Sep 17 00:00:00 2001 From: igneum-labs <337424239+igneum-labs@users.noreply.github.com> Date: Mon, 5 Oct 2026 20:49:45 +0000 Subject: [PATCH 3/3] read-width: per-watt and per-pound rows (telemetry watts, list prices approximate) and the AMD consequence (C11); packbench prints Metal's currentAllocatedSize and recommendedMaxWorkingSetSize (C12) Co-Authored-By: Claude Fable 5.1 --- docs/plans/read-width.md | 17 +++++++++++++++++ proto-metal/packbench.swift | 5 ++++- 2 files changed, 21 insertions(+), 1 deletion(-) diff --git a/docs/plans/read-width.md b/docs/plans/read-width.md index 60a4d3da8..b9c04b2e5 100644 --- a/docs/plans/read-width.md +++ b/docs/plans/read-width.md @@ -51,6 +51,23 @@ Probe ceilings at 1024 MiB (G dependent reads/s): RTX 5090 4 B 17.5, 16 B 18.0, | mix50-35-15 (6 programs, min / median / max) | 1,664 to 3,680 | 99.5 / 114.2 / 121.0, spread 18.8% | 17.45 / 18.76 / 18.83, 7.4% | 25.36 / 27.26 / 28.43, 11.3% | 6.1x | 5,939 / 8,192 expected | 0.620 | | mix25-50-25 (6 programs) | 2,240 to 5,024 | 95.9 / 107.3 / 119.8, 22.3% | 17.84 / 18.45 / 18.85, 5.5% | 23.21 / 24.68 / 25.21, 8.1% | 5.8x | 7,168 / 8,192 expected | 0.614 | +### 4.1 Per watt and per pound (consequences review C11) + +The runs carried no power sampling; the watts are the telemetry entry's (`docs/bench-log.md`, opencl-rdna4-telemetry, 5 October 2026: the RTX 5090 at 307.6 W under its 450 W cap for 122.3 MH/s, the RX 9070 XT at 199 W of its 304 W rating for about 17.8 MH/s, both on v2 with the shader clock at its top and the die waiting on memory), held constant across classes because every class is memory-bound on both cards (approximate: a class that moves more bytes per hash draws somewhat more at the memory controller, unmeasured). The Mac's GPU power is not measurable without root (`powermetrics`) and is taken as about 50 W (approximate, from memory). Prices are UK list, approximate, from memory. + +| Class | RTX 5090 MH/W (at 307.6 W) | RX 9070 XT MH/W (at 199 W) | 5090 / 9070 per watt | M5 Max MH/W (at about 50 W GPU, approximate) | 5090 MH per pound (at about 1,900, approximate) | 9070 XT MH per pound (at about 570, approximate) | +|---|---|---|---|---|---|---| +| v2 (w4) | 0.442 | 0.091 | 4.9x | 0.55 | 0.072 | 0.032 | +| w16 | 0.454 | 0.090 | 5.1x | 0.57 | 0.074 | 0.031 | +| w64 | 0.234 | 0.088 | 2.6x | 0.57 | 0.038 | 0.031 | +| w64x4 | 0.895 | 0.378 | 2.4x | 2.19 | 0.145 | 0.132 | +| mix50-35-15 (median) | 0.371 | 0.094 | 3.9x | 0.55 | 0.060 | 0.033 | +| mix25-50-25 (median) | 0.349 | 0.093 | 3.8x | 0.49 | 0.056 | 0.032 | +| scr8k32 | 0.397 | 0.071 | 5.6x | 0.98 | 0.064 | 0.025 | +| scr2k32 | 0.372 | 0.074 | 5.1x | 0.52 | 0.060 | 0.026 | + +Reading: whatever width is chosen, an AMD home miner keeps about a seventh of a 5090's rate and pays about 4.5x the electricity per hash, because every width costs the 9070 XT the same 2.4 G line fetches a second; per pound of card the 5090 is 2.2x the 9070 XT at v2 and w16 (0.072 against 0.032 MH/s per pound) and 4.9x per watt; only w64x4 narrows the per-pound gap (0.145 against 0.132), and that class fails the width rule. The consequence for the decision (D6, the project lead's): AMD's line width is not a read-width question at all; it is the card's random-access rate, and the levers that act on it (the 64 MB Infinity Cache against the dataset size, the memory path) are v3-or-3.0 questions outside this experiment. + Scratch, variant 5 (N persistent warps; GPU cost against the persistent control scr0k32; working set = 1 GiB + 256 MiB + 128 MiB output + N x size): | Class | RMW share | dataset B/hash | scratch B/hash read + written | RTX 5090 MH/s, 2,048 warps of 4,080 resident (vs control, share) | M5 Max Metal (vs control) | RX 9070 XT, 4,096 warps | working set 5090 / 9070 / Mac | diff --git a/proto-metal/packbench.swift b/proto-metal/packbench.swift index 3110bbe12..c708f5c28 100644 --- a/proto-metal/packbench.swift +++ b/proto-metal/packbench.swift @@ -204,6 +204,9 @@ print("pack \(packName) seed \"\(seedString)\" id \(programId) class \(className print("device \(device.name); compile \(String(format: "%.0f", compileMs)) ms; cache fill \(String(format: "%.1f", cacheGpu)) ms GPU (\(String(format: "%.1f", cacheWall)) wall); dataset build \(String(format: "%.1f", buildGpu)) ms GPU (\(String(format: "%.1f", buildWall)) wall)") print("cache FNV-1a 64 \(String(format: "%016llx", cacheFnv)) \(cacheOk ? "PASS" : "FAIL"); dataset head and last \(dsOk ? "PASS" : "FAIL"); vectors standalone \(vecPass)/\(vecBases.count), in batch \(batchVecPass)/\(batchVecN)") print("warm-up batch \(nonces) hashes: \(String(format: "%.1f", warmGpu)) ms GPU, \(String(format: "%.1f", warmWall)) ms wall") +let footprintMiB = device.currentAllocatedSize / 1048576 +let recommendedMiB = device.recommendedMaxWorkingSetSize / 1048576 +print("resident footprint after the timed batches: currentAllocatedSize \(footprintMiB) MiB (recommendedMaxWorkingSetSize \(recommendedMiB) MiB)") let overall = cacheOk && dsOk && vecPass == vecBases.count && batchVecPass == batchVecN -print("RESULT pack=\(packName) class=\(className) device=\(device.name.replacingOccurrences(of: " ", with: "_")) group=\(opts.group) warps=\(warpsN) arena_mib=\(persistent ? warpsN * 32 * scratchWordsPerLane * 4 / 1048576 : 0) nonces=\(nonces) batches=\(opts.batches) vectors=\(vecPass)/\(vecBases.count) batch_vectors=\(batchVecPass)/\(batchVecN) cache=\(cacheOk ? "PASS" : "FAIL") dataset=\(dsOk ? "PASS" : "FAIL") fingerprint=\(String(format: "%016llx", fingerprint)) mhs_gpu=\(String(format: "%.3f", mhsGpu)) mhs_wall=\(String(format: "%.3f", mhsWall)) loads=\(loadsPerHash) bytes=\(bytesPerHash) scratch_ops=\(scratchOps * 8) overall=\(overall ? "PASS" : "FAIL")") +print("RESULT pack=\(packName) class=\(className) device=\(device.name.replacingOccurrences(of: " ", with: "_")) group=\(opts.group) warps=\(warpsN) arena_mib=\(persistent ? warpsN * 32 * scratchWordsPerLane * 4 / 1048576 : 0) nonces=\(nonces) batches=\(opts.batches) vectors=\(vecPass)/\(vecBases.count) batch_vectors=\(batchVecPass)/\(batchVecN) cache=\(cacheOk ? "PASS" : "FAIL") dataset=\(dsOk ? "PASS" : "FAIL") fingerprint=\(String(format: "%016llx", fingerprint)) mhs_gpu=\(String(format: "%.3f", mhsGpu)) mhs_wall=\(String(format: "%.3f", mhsWall)) loads=\(loadsPerHash) bytes=\(bytesPerHash) footprint_mib=\(footprintMiB) recommended_mib=\(recommendedMiB) scratch_ops=\(scratchOps * 8) overall=\(overall ? "PASS" : "FAIL")") exit(overall ? 0 : 1)