Merge pull request 'readme: 4090 at 727 MH/s on the loop kernel, how to select it' (#30) from docs/readme-lair2 into main

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Ue5ZZm1Hiv5zPnucykKKuF

.gitea/workflows/deploy.yaml

@@ -20,6 +20,7 @@ on:
       - "**.md"
       - .gitea/workflows/ci.yml
       - .gitea/workflows/bench.yaml
+      - .gitea/workflows/release.yaml
   workflow_dispatch:
     inputs:
       mode:
@@ -87,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
@@ -145,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"
@@ -209,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 \

.gitea/workflows/release.yaml

@@ -0,0 +1,78 @@
+---
+# Release binaries for independent miners (quantus/miner#27).
+#
+# Pushing a tag `v<upstream version>-lair.<n>` builds the miner and the bench
+# harness on the cuda-13.0 runner (Fedora 43 like the mining hosts, and the
+# only runner with nvcc, so the binary carries the sm_86/sm_89/sm_120 cubins
+# and PTX), records what was built, and publishes a Gitea release with the
+# binaries, their checksums and the build report attached. The tag's commit
+# is embedded in `--version` so a downloaded binary identifies itself.
+#
+# Deploying to the fleet is deploy.yaml's job and follows main; a release is
+# only a public snapshot and changes nothing on the hosts.
+name: release
+
+on:
+  push:
+    tags: ["v*-lair.*"]
+
+env:
+  CARGO_TERM_COLOR: always
+
+jobs:
+  build:
+    runs-on: cuda-13.0
+    steps:
+      - uses: actions/checkout@v4
+      - name: build quantus-miner and quantus-bench
+        env:
+          MINER_BUILD_SHA: ${{ github.sha }}
+          MINER_CUDA_REQUIRE: "1"
+        run: |
+          set -euo pipefail
+          cargo build --release --locked -p miner-cli -p bench-harness
+          ./target/release/quantus-miner --version
+          ./target/release/quantus-bench --version || true
+      - name: package
+        run: |
+          set -euo pipefail
+          tag="${GITHUB_REF_NAME}"
+          dir="quantus-miner-${tag}-linux-x86_64"
+          mkdir -p "dist/${dir}"
+          cp target/release/quantus-miner target/release/quantus-bench "dist/${dir}/"
+          cp LICENSE "dist/${dir}/"
+          {
+            echo "quantus-miner ${tag}"
+            echo "commit: ${GITHUB_SHA}"
+            echo "built: $(date -u +%Y-%m-%dT%H:%M:%SZ) on $(source /etc/os-release && echo "$PRETTY_NAME"), $(ldd --version | head -1)"
+            echo "cuda: $(nvcc --version | grep -oE 'release [0-9.]+' | head -1)"
+            echo "cubins: sm_86 sm_89 sm_120 + compute_120 PTX"
+            echo "version string: $(./target/release/quantus-miner --version)"
+          } > "dist/${dir}/BUILD.txt"
+          (cd dist && tar -czf "${dir}.tar.gz" "${dir}" && sha256sum "${dir}.tar.gz" > "${dir}.tar.gz.sha256")
+          cat "dist/${dir}/BUILD.txt" "dist/${dir}.tar.gz.sha256"
+      - name: publish release
+        run: |
+          set -euo pipefail
+          tag="${GITHUB_REF_NAME}"
+          dir="quantus-miner-${tag}-linux-x86_64"
+          api="${GITHUB_API_URL}/repos/${GITHUB_REPOSITORY}"
+          auth="Authorization: token ${{ secrets.GITHUB_TOKEN }}"
+          body=$(python3 - "$tag" "dist/${dir}/BUILD.txt" "dist/${dir}.tar.gz.sha256" <<'PY'
+          import json, sys, pathlib
+          tag, build, sha = sys.argv[1], pathlib.Path(sys.argv[2]).read_text(), pathlib.Path(sys.argv[3]).read_text().strip()
+          text = (
+              f"quantus-miner {tag}: Linux x86_64, NVIDIA RTX 30/40/50 (sm_86, sm_89, sm_120, plus PTX), "
+              "glibc 2.42 or newer, driver with CUDA 13.0 support.\n\n"
+              "No dev fee, no licence server, no telemetry. Measured rates and how to reproduce them are in the README.\n\n"
+              f"```\n{build}```\n\nsha256: `{sha}`\n"
+          )
+          print(json.dumps({"tag_name": tag, "name": f"quantus-miner {tag}", "body": text, "draft": False, "prerelease": False}))
+          PY
+          )
+          id=$(curl -fsS -X POST -H "$auth" -H "Content-Type: application/json" "${api}/releases" -d "$body" | python3 -c 'import json,sys; print(json.load(sys.stdin)["id"])')
+          echo "release id ${id}"
+          for f in "dist/${dir}.tar.gz" "dist/${dir}.tar.gz.sha256" "dist/${dir}/BUILD.txt"; do
+            curl -fsS -X POST -H "$auth" "${api}/releases/${id}/assets?name=$(basename "$f")" -F "attachment=@${f}" > /dev/null
+            echo "attached $(basename "$f")"
+          done

README.md

@@ -2,6 +2,58 @@
 
 High-performance external mining service for Quantus Network with support for CPU, GPU, and hybrid CPU+GPU mining.
 
+## This fork
+
+Like the official miner, but fast: a native CUDA engine for NVIDIA cards, kept
+bit-exact with the reference hash and measured against every other miner we
+can get hold of. No dev fee, no licence server, no telemetry, solo or pool.
+Apache-2.0, like upstream.
+
+**Measured hashrate**, three interleaved 30-second rounds per card at the
+enforced power cap, memory clock locked at 810 MHz, CUDA 13.0 driver 580:
+
+| card | power cap | this fork | upstream `--cuda-gpu` | quantusminer.com qpow-cuda 1.0.7 | quanpool-miner 6.2.4 (5% fee) |
+|---|---|---|---|---|---|
+| RTX 5090 | 400 W | **1104 MH/s** | 944 | 1171 | 1234 (1172 after fee) |
+| RTX 4090 | 250 W | **727 MH/s** (loop kernel) | | | |
+| RTX 3060 | 130 W | **141 MH/s** | | | |
+
+The 5090 column was measured on the same card in the same session
+(2026-09-13); the closed pool binaries were run in benchmark mode only. How
+the pool kernels get their remaining edge, and what has been tried, is
+written up in [issue #27](https://git.lair.cafe/quantus/miner/issues/27).
+
+Since v4.0.2-lair.2 the binary carries two kernels. The default (`unrolled`)
+is fastest on the RTX 5090 and 3060; RTX 40-series cards measure 1.7% better
+on the rolled-loop one, selected with `MINER_CUDA_KERNEL=loop` in the
+environment (it reports as kernel `cuda-loop` on the metrics). Both are
+bit-exact with the CPU reference; try each with `quantus-bench` on your card.
+
+**Binaries** are on the [releases page](https://git.lair.cafe/quantus/miner/releases):
+one Linux x86_64 `quantus-miner` carrying cubins for sm_86 (RTX 30), sm_89
+(RTX 40) and sm_120 (RTX 50) plus PTX for anything newer, built on Fedora 43
+(glibc 2.42 or newer), and `quantus-bench`, the measurement harness. Nothing
+is compiled at run time; the driver needs CUDA 13.0 support (580.x).
+
+**Reproduce the numbers.** `quantus-bench` pins the power limit it expects,
+warms up, reports the median of timed windows with their spread, and checks
+GPU hashes against the CPU reference for random jobs:
+
+```bash
+# hashrate: median of 5 x 30 s windows, one worker thread per card
+./quantus-bench --duration-secs 30 --runs 5 --workers 1 --expect-power-limit 400
+
+# parity: 2000 random jobs, every found hash recomputed on the CPU
+./quantus-bench --duration-secs 3 --runs 1 --parity-jobs 2000
+```
+
+Compare builds with interleaved rounds (A, B, A, B, ...) rather than one
+block of runs each: on a power-capped card, clock drift between blocks looks
+like a 0.5% code change.
+
+**Solo mining** against your own node is the `serve` command below with the
+node's auth token and certificate fingerprint; nothing else is needed.
+
 ## Building
 
 ```bash

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}}