Merge pull request 'engine-cuda: loop-structured mining_loop kernel, opt-in via MINER_CUDA_KERNEL=loop' (#29) from perf/loop-kernel into main

.gitea/workflows/deploy.yaml

@@ -88,17 +88,24 @@ jobs:
           # `kernel` is the kernel id validate expects on the per-device metric:
           # cuda (engine-cuda, #3) on every NVIDIA host once the build carries
           # the fat binary; u64 would mean the miner silently fell back to wgpu.
+          # `cuda_kernel` selects the kernel entry (MINER_CUDA_KERNEL, #29):
+          # unrolled is the default; loop measured +1.7% on the 4090 and
+          # -1.8% / -3.4% on the 3060 / 5090 (2026-09-14), so only benjy runs
+          # it. `kernel` must be the id that choice reports: cuda or cuda-loop.
           - host: benjy.hanzalova.internal
             node: bob.hanzalova.internal
             gpu_devices: "1"        # 1x RTX 4090 (sm_89); reference card for #2
-            kernel: cuda
+            cuda_kernel: loop
+            kernel: cuda-loop
           - host: quadbrat.hanzalova.internal
             node: bob.hanzalova.internal
             gpu_devices: "1"        # 1x RTX 3060 (sm_86)
+            cuda_kernel: unrolled
             kernel: cuda
           - host: beast.hanzalova.internal
             node: bob.hanzalova.internal
             gpu_devices: "2"        # 2x RTX 5090 (sm_120)
+            cuda_kernel: unrolled
             kernel: cuda
     steps:
       - uses: actions/checkout@v4
@@ -146,6 +153,7 @@ jobs:
         if: ${{ github.event.inputs.mode != 'validate' }}
         env:
           GPU_DEVICES: ${{ matrix.gpu_devices }}
+          CUDA_KERNEL: ${{ matrix.cuda_kernel }}
         run: |
           set -euo pipefail
           SSHOPTS="-i $HOME/.ssh/id_gitea_ci -o StrictHostKeyChecking=accept-new"
@@ -210,6 +218,7 @@ jobs:
           t = pathlib.Path("deploy/miner.env.tmpl").read_text()
           t = t.replace("{{QUANTUS_NODE_ADDR}}", os.environ["NODE_ADDR"])
           t = t.replace("{{QUANTUS_GPU_DEVICES}}", os.environ["GPU_DEVICES"])
+          t = t.replace("{{MINER_CUDA_KERNEL}}", os.environ["CUDA_KERNEL"])
           pathlib.Path("miner.env").write_text(t)
           PY
           push --chown=root:quantus-miner --chmod=F0640 \

crates/engine-cuda/src/kernels/mining.cu

@@ -333,6 +333,32 @@ __device__ __forceinline__ u64 gf_canon(u64 a) {
 #else
 #define LAIR_LAUNCH_BOUNDS __launch_bounds__(LAIR_TPB)
 #endif
+// lair: the loop-structured kernel (`mining_loop`, quantus/miner#27). Same
+// arithmetic and same host contract as `mining_main`, but the round loops
+// are rolled with the constants fetched from constant memory by index, the
+// way the fastest closed kernel is built, so the register footprint can drop
+// and more warps fit per SM. Selected at run time (MINER_CUDA_KERNEL=loop);
+// the unrolled kernel stays the default until this one measures better.
+// 512 measured best for this kernel on beast (2131 MH/s against 2095 at
+// 256 and 2099 at 128 x 4 blocks); the host picks the same block size when
+// the loop kernel is selected.
+#ifndef LAIR_LOOP_TPB
+#define LAIR_LOOP_TPB 512
+#endif
+#ifndef LAIR_LOOP_MINBLOCKS
+#define LAIR_LOOP_MINBLOCKS 0
+#endif
+#if LAIR_LOOP_MINBLOCKS
+#define LAIR_LOOP_LAUNCH_BOUNDS __launch_bounds__(LAIR_LOOP_TPB, LAIR_LOOP_MINBLOCKS)
+#else
+#define LAIR_LOOP_LAUNCH_BOUNDS __launch_bounds__(LAIR_LOOP_TPB)
+#endif
+// Internal rounds per loop iteration in the loop kernel (1 or 2). One
+// measured 2105 MH/s against 1976 for two on beast: the pair doubles the
+// live state across the iteration for no scheduling gain.
+#ifndef LAIR_LOOP_INT_PAIR
+#define LAIR_LOOP_INT_PAIR 1
+#endif
 // Decide most nonces on element 0 of the final state before computing the
 // rest of the last linear layer: at mainnet difficulty the top ~47 bits of
 // the hash must be zero, and they live in that element.
@@ -708,6 +734,65 @@ __device__ __forceinline__ void permute(u64 st[12], bool skip_last_ext = false,
 #endif
 }
 
+// lair: rolled-loop twin of permute for the loop kernel. Round-constant
+// rows are addressed by the loop variable, so they are constant-memory loads
+// with a uniform offset instead of immediates; the per-element loops inside
+// each round stay unrolled so the state lives in registers.
+#if LAIR_LAZY_ADD && LAIR_FOLD_RC
+static_assert(N_INTERNAL % 2 == 0, "the paired internal loop assumes an even round count");
+__device__ __forceinline__ void permute_loop(u64 st[12], bool skip_last_ext, bool skip_first_ext) {
+    if (!skip_first_ext) ext_layer_rc(st, RC_INITIAL[0]);
+#pragma unroll 1
+    for (int r = 0; r + 1 < N_EXTERNAL_HALF; r++) {
+#pragma unroll
+        for (int i = 0; i < 12; i++) st[i] = gf_sbox(st[i]);
+        ext_layer_rc(st, RC_INITIAL[r + 1]);
+    }
+#pragma unroll
+    for (int i = 0; i < 12; i++) st[i] = gf_sbox(st[i]);
+    ext_layer(st);
+    st[0] = gf_add(st[0], RC_INTERNAL[0]);
+#if LAIR_LOOP_INT_PAIR == 2
+#pragma unroll 1
+    for (int r = 0; r + 2 < N_INTERNAL; r += 2) {
+        st[0] = gf_sbox(st[0]);
+        int_layer_rc(st, RC_INTERNAL[r + 1], true, nullptr);
+        st[0] = gf_sbox(st[0]);
+        int_layer_rc(st, RC_INTERNAL[r + 2], true, nullptr);
+    }
+    st[0] = gf_sbox(st[0]);
+    int_layer_rc(st, RC_INTERNAL[N_INTERNAL - 1], true, nullptr);
+#else
+#pragma unroll 1
+    for (int r = 0; r + 1 < N_INTERNAL; r++) {
+        st[0] = gf_sbox(st[0]);
+        int_layer_rc(st, RC_INTERNAL[r + 1], true, nullptr);
+    }
+#endif
+    st[0] = gf_sbox(st[0]);
+    int_layer_rc(st, 0ull, false, RC_TERMINAL[0]);
+#pragma unroll 1
+    for (int r = 0; r + 1 < N_EXTERNAL_HALF; r++) {
+#pragma unroll
+        for (int i = 0; i < 12; i++) st[i] = gf_sbox(st[i]);
+        ext_layer_rc(st, RC_TERMINAL[r + 1]);
+    }
+#pragma unroll
+    for (int i = 0; i < 12; i++) st[i] = gf_sbox(st[i]);
+    if (!skip_last_ext) ext_layer(st);
+}
+#else
+__device__ __forceinline__ void permute_loop(u64 st[12], bool skip_last_ext, bool skip_first_ext) {
+    permute(st, skip_last_ext, skip_first_ext);
+}
+#endif
+
+template <bool LOOP>
+__device__ __forceinline__ void run_permute(u64 st[12], bool skip_last_ext = false, bool skip_first_ext = false) {
+    if (LOOP) permute_loop(st, skip_last_ext, skip_first_ext);
+    else permute(st, skip_last_ext, skip_first_ext);
+}
+
 __device__ __forceinline__ u32 bswap32(u32 v) {
     return __byte_perm(v, 0, 0x0123);
 }
@@ -742,12 +827,14 @@ static_assert(sizeof(MiningUniforms) == 320, "MiningUniforms must match the host
 __device__ __constant__ u64 LAIR_DIR7[12] = {1, 1, 3, 2, 2, 2, 6, 4, 1, 1, 3, 2};
 #endif
 
-extern "C" __global__ void LAIR_LAUNCH_BOUNDS
-mining_main(u32* __restrict__ results,
-            const MiningUniforms uni,
-            u32 total_threads,
-            u32 nonces_per_thread,
-            u32 total_nonces)
+// The per-nonce search, shared by both kernel entries; LOOP picks the
+// permutation implementation at compile time.
+template <bool LOOP>
+__device__ __forceinline__ void mine_nonces(u32* __restrict__ results,
+                                            const MiningUniforms& uni,
+                                            u32 total_threads,
+                                            u32 nonces_per_thread,
+                                            u32 total_nonces)
 {
     if (*((volatile u32*)results) != 0u) return;
     u32 tid = blockIdx.x * blockDim.x + threadIdx.x;
@@ -796,13 +883,13 @@ mining_main(u32* __restrict__ results,
             for (int i = 0; i < 8; i++) st[i] = gf_add(st[i], (u64)bswap32(current_nonce[7 - i]));
             ext_layer_rc(st, RC_INITIAL[0]);
         }
-        permute(st, false, true);
+        run_permute<LOOP>(st, false, true);
 #else
 #pragma unroll
         for (int i = 0; i < 12; i++) st[i] = uni.midstate[i];
 #pragma unroll
         for (int i = 0; i < 8; i++) st[i] = gf_add(st[i], (u64)bswap32(current_nonce[7 - i]));
-        permute(st);
+        run_permute<LOOP>(st);
 #endif
         st[0] = gf_add(st[0], 1ull);
         st[1] = gf_add(st[1], 1ull);
@@ -810,7 +897,7 @@ mining_main(u32* __restrict__ results,
         // Element 0 of the final state is the hash's top 64 bits. Compute it
         // alone, and only nonces that do not already exceed the target there
         // pay for the other eleven outputs of the last linear layer.
-        permute(st, true);
+        run_permute<LOOP>(st, true);
         {
             u64 c0 = gf_canon(ext_layer_out0(st));
             u32 h0 = bswap32((u32)(c0 & EPS));
@@ -820,7 +907,7 @@ mining_main(u32* __restrict__ results,
         }
         ext_layer(st);
 #else
-        permute(st);
+        run_permute<LOOP>(st);
 #endif
 
         // First squeeze: most significant 256 bits of the hash.
@@ -843,7 +930,7 @@ mining_main(u32* __restrict__ results,
         u32 hash_le[16];
 #pragma unroll
         for (int i = 0; i < 8; i++) hash_le[15 - i] = bswap32(first[i]);
-        permute(st);
+        run_permute<LOOP>(st);
 #pragma unroll
         for (int i = 0; i < 4; i++) {
             u64 c = gf_canon(st[i]);
@@ -874,3 +961,23 @@ mining_main(u32* __restrict__ results,
         }
     }
 }
+
+extern "C" __global__ void LAIR_LAUNCH_BOUNDS
+mining_main(u32* __restrict__ results,
+            const MiningUniforms uni,
+            u32 total_threads,
+            u32 nonces_per_thread,
+            u32 total_nonces)
+{
+    mine_nonces<false>(results, uni, total_threads, nonces_per_thread, total_nonces);
+}
+
+extern "C" __global__ void LAIR_LOOP_LAUNCH_BOUNDS
+mining_loop(u32* __restrict__ results,
+            const MiningUniforms uni,
+            u32 total_threads,
+            u32 nonces_per_thread,
+            u32 total_nonces)
+{
+    mine_nonces<true>(results, uni, total_threads, nonces_per_thread, total_nonces);
+}

crates/engine-cuda/src/lib.rs

@@ -26,18 +26,60 @@ use std::time::Instant;
 const FATBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/mining.fatbin"));
 /// SM list the fat binary carries, for `--version` and the config metric.
 pub const CUDA_ARCHS: &str = env!("MINER_CUDA_ARCHS");
-const KERNEL_ID: &str = "cuda";
+/// lair: which kernel entry the module runs (quantus/miner#27). `unrolled`
+/// is `mining_main`, the deployed kernel; `loop` is `mining_loop`, the
+/// rolled-round twin under evaluation. Selected with MINER_CUDA_KERNEL so a
+/// host can be switched without a rebuild, and reported as the kernel id on
+/// the per-device metrics so the deploy validate and the dashboards can tell
+/// them apart.
+#[derive(Clone, Copy, PartialEq, Eq, Debug)]
+enum KernelChoice {
+    Unrolled,
+    Loop,
+}
+
+impl KernelChoice {
+    fn from_env() -> Self {
+        match std::env::var("MINER_CUDA_KERNEL").as_deref() {
+            Ok("loop") => KernelChoice::Loop,
+            Ok("unrolled") | Ok("") | Err(_) => KernelChoice::Unrolled,
+            Ok(other) => {
+                log::warn!(target: "cuda_engine", "MINER_CUDA_KERNEL={other:?} is not unrolled|loop; using unrolled");
+                KernelChoice::Unrolled
+            }
+        }
+    }
+    fn function(self) -> &'static str {
+        match self {
+            KernelChoice::Unrolled => "mining_main",
+            KernelChoice::Loop => "mining_loop",
+        }
+    }
+    /// Kernel id on metrics: `cuda` for the deployed kernel (what the deploy
+    /// validate asserts), `cuda-loop` for the candidate.
+    fn id(self) -> &'static str {
+        match self {
+            KernelChoice::Unrolled => "cuda",
+            KernelChoice::Loop => "cuda-loop",
+        }
+    }
+}
 /// Threads per block. Must match the kernel's `__launch_bounds__` (LAIR_TPB,
 /// default 256); override for experiments with MINER_CUDA_THREADS_PER_BLOCK
 /// on a kernel compiled with the same value.
 const THREADS_PER_BLOCK_DEFAULT: u32 = 256;
+/// The loop kernel's `LAIR_LOOP_TPB`.
+const LOOP_THREADS_PER_BLOCK_DEFAULT: u32 = 512;
 
-fn threads_per_block() -> u32 {
+fn threads_per_block(kernel: KernelChoice) -> u32 {
     std::env::var("MINER_CUDA_THREADS_PER_BLOCK")
         .ok()
         .and_then(|v| v.parse().ok())
         .filter(|v: &u32| *v > 0 && v.is_multiple_of(32))
-        .unwrap_or(THREADS_PER_BLOCK_DEFAULT)
+        .unwrap_or(match kernel {
+            KernelChoice::Unrolled => THREADS_PER_BLOCK_DEFAULT,
+            KernelChoice::Loop => LOOP_THREADS_PER_BLOCK_DEFAULT,
+        })
 }
 
 /// lair: size every batch to whole grids. The CLI default batch (1M nonces)
@@ -93,6 +135,7 @@ struct Device {
     func: CudaFunction,
     name: String,
     sm_count: u32,
+    kernel: KernelChoice,
     threads_per_block: u32,
     threads_per_sm: u32,
     metrics: metrics::DeviceMetrics,
@@ -171,9 +214,10 @@ impl CudaEngine {
                 Ok(d) => {
                     log::info!(
                         target: "cuda_engine",
-                        "CUDA device {ordinal}: {} ({} SMs) using {KERNEL_ID} kernel [sm_{{{}}}]",
+                        "CUDA device {ordinal}: {} ({} SMs) using {} kernel [sm_{{{}}}]",
                         d.name,
                         d.sm_count,
+                        d.kernel.id(),
                         CUDA_ARCHS
                     );
                     devices.push(Arc::new(d));
@@ -221,15 +265,17 @@ impl Device {
             cudarc::driver::sys::CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT,
         )? as u32;
         let module = ctx.load_module(Ptx::from_binary(FATBIN.to_vec()))?;
-        let func = module.load_function("mining_main")?;
+        let kernel = KernelChoice::from_env();
+        let func = module.load_function(kernel.function())?;
         Ok(Self {
             ctx,
             func,
             name,
             sm_count,
+            kernel,
             threads_per_sm: threads_per_sm(),
-            threads_per_block: threads_per_block(),
-            metrics: metrics::DeviceMetrics::new(ordinal, KERNEL_ID),
+            threads_per_block: threads_per_block(kernel),
+            metrics: metrics::DeviceMetrics::new(ordinal, kernel.id()),
         })
     }
 

deploy/miner.env.tmpl

@@ -4,3 +4,5 @@
 # from the node host as separate 0640 files.
 QUANTUS_NODE_ADDR={{QUANTUS_NODE_ADDR}}
 QUANTUS_GPU_DEVICES={{QUANTUS_GPU_DEVICES}}
+# Kernel entry in the CUDA engine: unrolled (mining_main) or loop (mining_loop), quantus/miner#29.
+MINER_CUDA_KERNEL={{MINER_CUDA_KERNEL}}