Merge pull request 'deploy: tie the GPU profile to the miner's lifecycle' (#33) from feat/gpu-profile-with-lifecycle into main

The power limit and 810 MHz memory lock were applied at boot as a host
setting, so they outlived the miner: a host switched back to LLM
inference kept the lock and decoded at a sixth of its measured speed,
with nothing reporting as throttled.

- quantus-miner.service applies the profile in ExecStartPre and resets
  to driver defaults in ExecStopPost (nvidia-power-limit.sh, installed
  by lair/quantus). ExecStopPost runs on stop, crash and failed start.
- deploy: on a host with neuron.service enabled, ship the files but do
  not enable or start the miner; validate passes there without checking
  hashing. neuron conflicts with the miner, so starting it would stop
  inference behind the operator's back.
- bench-on-host.sh: stopping the unit now resets the clocks, so apply
  the miner profile for the measurement window (and read the asserted
  power limit after it), restore the prior profile afterwards, and
  refuse to measure on a host where neuron is active.
- infra-setup.sh: sudoers for the two script invocations the bench
  needs. Re-run it on the mining hosts before the next deploy; the
  preflight will name the missing path until then.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01XViS3RVfGTRaRSopLN68MA

.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:
@@ -60,6 +61,11 @@ jobs:
           # Embedded by crates/miner-cli/build.rs into --version; validate below
           # asserts the host runs exactly this commit.
           MINER_BUILD_SHA: ${{ github.sha }}
+          # Fail the build here if nvcc is missing rather than shipping a
+          # binary that falls back to wgpu and failing validate on the hosts.
+          # (The same warning in ci.yml's test job is expected: the `rust`
+          # runner has no nvcc and only lints and tests.)
+          MINER_CUDA_REQUIRE: "1"
         run: |
           set -euo pipefail
           cargo build --release --locked -p miner-cli
@@ -82,22 +88,30 @@ jobs:
           # so the firewalld rich rule admits it. lair/quantus's SCRAPE_MINERS
           # must list it so Prometheus scrapes it.
           #
-          # Mainnet call 2026-09-09 (quantus/miner#1): all three hosts mine;
-          # neuron (cortex inference) is disabled on each of them.
+          # Mainnet call 2026-09-09 (quantus/miner#1): all three hosts mine
+          # when neuron (helexa inference) is not enabled on them. A host with
+          # neuron enabled keeps its miner files current but is not started.
           # `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 +159,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 +224,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 \
@@ -248,6 +264,20 @@ jobs:
             run "sudo firewall-cmd --zone=$zone --add-rich-rule='$mrich'"
           fi
 
+          # A host whose neuron.service is enabled is serving LLM inference.
+          # neuron conflicts with the miner (lair/quantus drop-in), so enabling
+          # or starting the miner here would stop inference mid-request and put
+          # the host back to mining behind the operator's back. Ship the files,
+          # reload so the next deliberate start sees them, and leave it there.
+          # Switching a host is `systemctl disable --now neuron && systemctl
+          # enable --now quantus-miner`, done on purpose.
+          if run systemctl is-enabled --quiet neuron.service; then
+            run sudo systemctl daemon-reload
+            echo "neuron.service is enabled on ${{ matrix.host }}: inference host, miner left stopped"
+            echo "restarted=0" >> "$GITHUB_OUTPUT"
+            exit 0
+          fi
+
           run sudo systemctl enable quantus-miner.service   # idempotent
           if [ "$RESTART" = 1 ]; then
             echo "changes applied — restarting"
@@ -269,6 +299,14 @@ jobs:
           run() { ssh $SSHOPTS gitea_ci@"${{ matrix.host }}" "$@"; }
           fail=0
 
+          # Deploy leaves the miner stopped on an inference host (see above);
+          # there is nothing to validate there, and failing would make every
+          # miner deploy red for as long as the host serves.
+          if run systemctl is-enabled --quiet neuron.service; then
+            echo "neuron.service is enabled on ${{ matrix.host }}: inference host, miner not expected to run"
+            exit 0
+          fi
+
           echo "--- unit (${{ matrix.host }} -> ${{ matrix.node }}) ---"
           run systemctl is-active quantus-miner.service
 

.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/bench-harness/bench-on-host.sh

@@ -22,10 +22,15 @@ fi
 
 echo "--- host ---"
 hostname -f
-limit=$(nvidia-smi --query-gpu=power.limit --format=csv,noheader,nounits | awk 'NR==1{printf "%d", $1}')
-echo "power limit: ${limit} W"
 "$BIN" --version
 
+# A host serving inference is not a bench host, even when idle: the window
+# locks the memory clock, which would cripple any request that lands in it.
+if systemctl is-active --quiet neuron.service; then
+  echo "neuron.service is active on this host; refusing to measure" >&2
+  exit 1
+fi
+
 echo "--- pause miner ---"
 was_active=0
 if systemctl is-active --quiet quantus-miner.service; then
@@ -37,8 +42,12 @@ fi
 resume() {
   if [ "$was_active" = 1 ]; then
     echo "--- resume miner ---"
+    # Its ExecStartPre applies the miner profile again.
     sudo systemctl start quantus-miner.service
     systemctl is-active quantus-miner.service
+  else
+    echo "--- restore default GPU profile ---"
+    sudo /usr/local/bin/nvidia-power-limit.sh default
   fi
 }
 trap resume EXIT
@@ -53,6 +62,14 @@ if [ "$util" -gt 5 ]; then
   exit 1
 fi
 
+# Stopping the miner reset the card to its default profile (ExecStopPost).
+# Measure under the profile the miner actually runs with, and read the limit
+# the harness asserts only after applying it.
+echo "--- apply miner GPU profile ---"
+sudo /usr/local/bin/nvidia-power-limit.sh miner
+limit=$(nvidia-smi --query-gpu=power.limit --format=csv,noheader,nounits | awk 'NR==1{printf "%d", $1}')
+echo "power limit: ${limit} W"
+
 echo "--- measure ---"
 cd "$DIR"
 RUST_LOG=info "$BIN" \

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

@@ -333,6 +333,47 @@ __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
+// Loop kernel: take the diagonal from the launch parameters instead of the
+// __constant__ table (see MiningUniforms::diag).
+#ifndef LAIR_LOOP_PARAM_DIAG
+#define LAIR_LOOP_PARAM_DIAG 1
+#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.
+//
+// Where the loop kernel's internal round (239 instructions at 128
+// registers) still exceeds the quanpool kernel's 199 (2026-09-14, static
+// screen on sm_120): 12 reloads of the diagonal constants per round, which
+// ptxas issues even with registers to spare (13 LDC at 192 registers too;
+// hoisting them in C or passing them as 32-bit halves compiles to identical
+// SASS); ~12 in the product carry assembly; ~8 in the EPS multiply of the
+// reductions; ~5 in the sum-plus-constant fold; ~6 moves. Five multiply-add
+// formulations (LAIR_MULADD) and two constant deliveries were screened and
+// none beat the shipped form.
+#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.
@@ -454,8 +495,15 @@ __device__ __forceinline__ u64 acc_reduce(Acc a) {
     return r;
 }
 
-// a * b + c mod P. The addend rides in the two `mad.wide` partial products
-// (a0*b0 + c_lo and a1*b0 + c_hi both fit in 64 bits), so it costs nothing.
+// a * b + c mod P. Formulations are screened by the loop kernel's internal
+// round body (static SASS count, target quanpool's 199) before any card
+// time: LAIR_MULADD selects one. 0 is the shipped form.
+#ifndef LAIR_MULADD
+#define LAIR_MULADD 0
+#endif
+#if LAIR_MULADD == 0
+// The addend rides in the two `mad.wide` partial products (a0*b0 + c_lo and
+// a1*b0 + c_hi both fit in 64 bits), so it costs nothing.
 __device__ __forceinline__ u64 gf_mul_add(u64 a, u64 b, u64 c) {
     u64 r;
     u32 e = LAIR_EPS32;
@@ -484,6 +532,120 @@ __device__ __forceinline__ u64 gf_mul_add(u64 a, u64 b, u64 c) {
         : "=l"(r) : "l"(a), "l"(b), "l"(c0), "l"(c1), "r"(e));
     return r;
 }
+#elif LAIR_MULADD == 1
+// Whole 64-bit addend folded into a0*b0 with a carry-out (mad.lo.cc /
+// madc.hi.cc); the middle products accumulated the same way. Measured worse
+// than 0 on 2026-09-14 (loop internal body 268 against 254).
+__device__ __forceinline__ u64 gf_mul_add(u64 a, u64 b, u64 c) {
+    u64 r;
+    u32 e = LAIR_EPS32;
+    asm("{\n\t"
+        ".reg .b32 a0, a1, b0, b1, s0, s1, m0, m1, cw, c2, ll, lh, hl, hh, c;\n\t"
+        ".reg .b64 m;\n\t"
+        "mov.b64 {a0, a1}, %1;\n\t"
+        "mov.b64 {b0, b1}, %2;\n\t"
+        "mov.b64 {s0, s1}, %3;\n\t"
+        "mul.wide.u32 m, a1, b0;\n\t"
+        "mov.b64 {m0, m1}, m;\n\t"
+        "mad.lo.cc.u32 m0, a0, b1, m0;\n\t"
+        "madc.hi.cc.u32 m1, a0, b1, m1;\n\t"
+        "addc.u32 cw, 0, 0;\n\t"
+        "mad.lo.cc.u32 ll, a0, b0, s0;\n\t"
+        "madc.hi.cc.u32 lh, a0, b0, s1;\n\t"
+        "addc.u32 c2, 0, 0;\n\t"
+        "add.cc.u32 lh, lh, m0;\n\t"
+        "addc.cc.u32 hl, m1, c2;\n\t"
+        "addc.u32 hh, cw, 0;\n\t"
+        "mad.lo.cc.u32 hl, a1, b1, hl;\n\t"
+        "madc.hi.u32 hh, a1, b1, hh;\n\t"
+        LAIR_REDUCE_PTX("%4")
+        "mov.b64 %0, {ll, lh};\n\t"
+        "}"
+        : "=l"(r) : "l"(a), "l"(b), "l"(c), "r"(e));
+    return r;
+}
+#elif LAIR_MULADD == 2
+// nvcc's own 128-bit product (a*b, __umul64hi), the addend folded in C, the
+// reduction as the PTX chain over C temporaries.
+__device__ __forceinline__ u64 gf_mul_add(u64 a, u64 b, u64 c) {
+    u64 lo = a * b + c;
+    u64 hi = __umul64hi(a, b) + (lo < c ? 1ull : 0ull);
+    u32 ll = (u32)lo, lh = (u32)(lo >> 32), hl = (u32)hi, hh = (u32)(hi >> 32), cc;
+    u32 e = LAIR_EPS32;
+    asm("mad.lo.cc.u32 %0, %3, %5, %0;\n\tmadc.hi.cc.u32 %1, %3, %5, %1;\n\taddc.u32 %2, %4, 0;\n\taddc.u32 %1, %1, 0;\n\tsub.cc.u32 %0, %0, %2;\n\tsubc.u32 %1, %1, 0;"
+        : "+r"(ll), "+r"(lh), "=r"(cc) : "r"(hl), "r"(hh), "r"(e));
+    return ((u64)lh << 32) | ll;
+}
+#elif LAIR_MULADD == 3
+// The shipped product chain (gf_mul's), then the addend added to the 128-bit
+// product as a 32-bit carry chain before the reduction.
+__device__ __forceinline__ u64 gf_mul_add(u64 a, u64 b, u64 c) {
+    u64 r;
+    u32 e = LAIR_EPS32;
+    asm("{\n\t"
+        ".reg .b32 a0, a1, b0, b1, s0, s1, p0l, p0h, m0, m1, cw, ll, lh, hl, hh, c;\n\t"
+        ".reg .b64 p0, m, m2, p3, t, hi;\n\t"
+        "mov.b64 {a0, a1}, %1;\n\t"
+        "mov.b64 {b0, b1}, %2;\n\t"
+        "mov.b64 {s0, s1}, %3;\n\t"
+        "mul.wide.u32 p0, a0, b0;\n\t"
+        "mul.wide.u32 m, a1, b0;\n\t"
+        "mul.wide.u32 m2, a0, b1;\n\t"
+        "add.cc.u64 m, m, m2;\n\t"
+        "addc.u32 cw, 0, 0;\n\t"
+        "mov.b64 {p0l, p0h}, p0;\n\t"
+        "mov.b64 {m0, m1}, m;\n\t"
+        "mov.b64 t, {m1, cw};\n\t"
+        "mul.wide.u32 p3, a1, b1;\n\t"
+        "add.cc.u32 lh, p0h, m0;\n\t"
+        "addc.u64 hi, p3, t;\n\t"
+        "mov.b64 {hl, hh}, hi;\n\t"
+        "add.cc.u32 ll, p0l, s0;\n\t"
+        "addc.cc.u32 lh, lh, s1;\n\t"
+        "addc.cc.u32 hl, hl, 0;\n\t"
+        "addc.u32 hh, hh, 0;\n\t"
+        LAIR_REDUCE_PTX("%4")
+        "mov.b64 %0, {ll, lh};\n\t"
+        "}"
+        : "=l"(r) : "l"(a), "l"(b), "l"(c), "r"(e));
+    return r;
+}
+#elif LAIR_MULADD == 4
+// Halves of the addend as mad.wide addends (as 0) but with the product's
+// carry assembly in 32-bit chains instead of add.cc.u64.
+__device__ __forceinline__ u64 gf_mul_add(u64 a, u64 b, u64 c) {
+    u64 r;
+    u32 e = LAIR_EPS32;
+    u64 c0 = c & EPS, c1 = c >> 32;
+    asm("{\n\t"
+        ".reg .b32 a0, a1, b0, b1, p0l, p0h, p3l, p3h, m0, m1, m2l, m2h, cw, ll, lh, hl, hh, c;\n\t"
+        ".reg .b64 p0, m, m2, p3;\n\t"
+        "mov.b64 {a0, a1}, %1;\n\t"
+        "mov.b64 {b0, b1}, %2;\n\t"
+        "mad.wide.u32 p0, a0, b0, %3;\n\t"
+        "mad.wide.u32 m, a1, b0, %4;\n\t"
+        "mul.wide.u32 m2, a0, b1;\n\t"
+        "mul.wide.u32 p3, a1, b1;\n\t"
+        "mov.b64 {p0l, p0h}, p0;\n\t"
+        "mov.b64 {p3l, p3h}, p3;\n\t"
+        "mov.b64 {m0, m1}, m;\n\t"
+        "mov.b64 {m2l, m2h}, m2;\n\t"
+        "add.cc.u32 m0, m0, m2l;\n\t"
+        "addc.cc.u32 m1, m1, m2h;\n\t"
+        "addc.u32 cw, p3h, 0;\n\t"
+        "add.cc.u32 lh, p0h, m0;\n\t"
+        "addc.cc.u32 hl, p3l, m1;\n\t"
+        "addc.u32 hh, cw, 0;\n\t"
+        "mov.b32 ll, p0l;\n\t"
+        LAIR_REDUCE_PTX("%5")
+        "mov.b64 %0, {ll, lh};\n\t"
+        "}"
+        : "=l"(r) : "l"(a), "l"(b), "l"(c0), "l"(c1), "r"(e));
+    return r;
+}
+#else
+#error "unknown LAIR_MULADD"
+#endif
 
 // a * b + (c.lo + c.hi * 2^64) mod P: reduce the addend first (4 ops), then
 // fold it into the product for free.
@@ -571,7 +733,12 @@ __device__ __forceinline__ u64 ext_layer_out0(const u64 st[12]) {
 // Internal linear layer: one deferred sum, folded into each diagonal multiply.
 // `rc0` is the next internal round's constant for element 0, folded into that
 // element's multiply-add; `rc_row` a full row (the first terminal round's).
+__device__ __forceinline__ void int_layer_rc_d(u64 st[12], u64 rc0, bool has_rc0, const u64* rc_row, const u64* d);
 __device__ __forceinline__ void int_layer_rc(u64 st[12], u64 rc0, bool has_rc0, const u64* rc_row) {
+    int_layer_rc_d(st, rc0, has_rc0, rc_row, MDS_DIAG);
+}
+#define LAIR_DMA(x, i, add) gf_mul_add((x), d[i], (add))
+__device__ __forceinline__ void int_layer_rc_d(u64 st[12], u64 rc0, bool has_rc0, const u64* rc_row, const u64* d) {
     Acc s = acc_of(st[0]);
 #pragma unroll
     for (int i = 1; i < 12; i++) s = acc_add(s, st[i]);
@@ -580,15 +747,15 @@ __device__ __forceinline__ void int_layer_rc(u64 st[12], u64 rc0, bool has_rc0,
         // The constant joins the unreduced sum: one more limb add, then the
         // same single reduction.
 #pragma unroll
-        for (int i = 0; i < 12; i++) st[i] = gf_mul_add_acc(st[i], MDS_DIAG[i], acc_add(s, rc_row[i]));
+        for (int i = 0; i < 12; i++) st[i] = LAIR_DMA(st[i], i, acc_reduce(acc_add(s, rc_row[i])));
     } else {
         if (has_rc0) {
-            st[0] = gf_mul_add_acc(st[0], MDS_DIAG[0], acc_add(s, rc0));
+            st[0] = LAIR_DMA(st[0], 0, acc_reduce(acc_add(s, rc0)));
         } else {
-            st[0] = gf_mul_add(st[0], MDS_DIAG[0], sum);
+            st[0] = LAIR_DMA(st[0], 0, sum);
         }
 #pragma unroll
-        for (int i = 1; i < 12; i++) st[i] = gf_mul_add(st[i], MDS_DIAG[i], sum);
+        for (int i = 1; i < 12; i++) st[i] = LAIR_DMA(st[i], i, sum);
     }
 }
 
@@ -708,6 +875,72 @@ __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, const u64* diag) {
+    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_PARAM_DIAG
+    const u64* d = diag;
+#else
+    const u64* d = MDS_DIAG;
+    (void)diag;
+#endif
+#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_d(st, RC_INTERNAL[r + 1], true, nullptr, d);
+        st[0] = gf_sbox(st[0]);
+        int_layer_rc_d(st, RC_INTERNAL[r + 2], true, nullptr, d);
+    }
+    st[0] = gf_sbox(st[0]);
+    int_layer_rc_d(st, RC_INTERNAL[N_INTERNAL - 1], true, nullptr, d);
+#else
+#pragma unroll 1
+    for (int r = 0; r + 1 < N_INTERNAL; r++) {
+        st[0] = gf_sbox(st[0]);
+        int_layer_rc_d(st, RC_INTERNAL[r + 1], true, nullptr, d);
+    }
+#endif
+    st[0] = gf_sbox(st[0]);
+    int_layer_rc_d(st, 0ull, false, RC_TERMINAL[0], d);
+#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, const u64* diag) {
+    (void)diag;
+    permute(st, skip_last_ext, skip_first_ext);
+}
+#endif
+
+template <bool LOOP>
+__device__ __forceinline__ void run_permute(u64 st[12], const u64* diag, bool skip_last_ext = false, bool skip_first_ext = false) {
+    if (LOOP) permute_loop(st, skip_last_ext, skip_first_ext, diag);
+    else permute(st, skip_last_ext, skip_first_ext);
+}
+
 __device__ __forceinline__ u32 bswap32(u32 v) {
     return __byte_perm(v, 0, 0x0123);
 }
@@ -732,8 +965,12 @@ struct MiningUniforms {
     // State after the first external layer + its round constant for the
     // batch's first nonce (host: first_layer_after_absorb).
     u64 layer0_base[12];
+    // The internal layer's diagonal, same values as MDS_DIAG, for the loop
+    // kernel: a parameter-bank operand costs no load and no register, where
+    // ptxas reloaded the __constant__ table every round.
+    u64 diag[12];
 };
-static_assert(sizeof(MiningUniforms) == 320, "MiningUniforms must match the host layout");
+static_assert(sizeof(MiningUniforms) == 416, "MiningUniforms must match the host layout");
 
 #if LAIR_NONCE_DIR
 // Column 7 of the external matrix circ(2*M4, M4, M4): what the layer adds per
@@ -742,12 +979,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 +1035,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, uni.diag, 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, uni.diag);
 #endif
         st[0] = gf_add(st[0], 1ull);
         st[1] = gf_add(st[1], 1ull);
@@ -810,7 +1049,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, uni.diag, true);
         {
             u64 c0 = gf_canon(ext_layer_out0(st));
             u32 h0 = bswap32((u32)(c0 & EPS));
@@ -820,7 +1059,7 @@ mining_main(u32* __restrict__ results,
         }
         ext_layer(st);
 #else
-        permute(st);
+        run_permute<LOOP>(st, uni.diag);
 #endif
 
         // First squeeze: most significant 256 bits of the hash.
@@ -843,7 +1082,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, uni.diag);
 #pragma unroll
         for (int i = 0; i < 4; i++) {
             u64 c = gf_canon(st[i]);
@@ -874,3 +1113,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,
@@ -130,9 +173,13 @@ struct MiningUniforms {
     /// derives each nonce's state from this by one scalar times a fixed
     /// small-integer column instead of recomputing the layer (quantus/miner#3).
     layer0_base: [u64; 12],
+    /// lair: the internal layer's diagonal, for the loop kernel (#27). In the
+    /// parameter bank ptxas reads it straight from the instruction operand;
+    /// from `__constant__` it reloaded all twelve values every round.
+    diag: [u64; 12],
 }
 
-// SAFETY: `#[repr(C)]` plain-old-data with no padding (12 x u64, 32 x u32, 12 x u64),
+// SAFETY: `#[repr(C)]` plain-old-data with no padding (12 x u64, 32 x u32, 24 x u64),
 // matching the kernel's parameter layout; there is nothing to validate beyond
 // the layout, which is what this marker asserts.
 unsafe impl DeviceRepr for MiningUniforms {}
@@ -140,7 +187,7 @@ unsafe impl DeviceRepr for MiningUniforms {}
 // A mismatch here would silently feed the kernel the wrong midstate or target,
 // so it is a build error rather than a mining bug. The kernel carries the same
 // assertion.
-const _: () = assert!(std::mem::size_of::<MiningUniforms>() == 320);
+const _: () = assert!(std::mem::size_of::<MiningUniforms>() == 416);
 
 pub struct CudaEngine {
     engine_id: usize,
@@ -171,9 +218,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 +269,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()),
         })
     }
 
@@ -493,6 +543,7 @@ impl CudaEngine {
             start_nonce: start_u32s,
             target: res.target_u32s,
             layer0_base,
+            diag: qp_poseidon_constants::POSEIDON2_MATRIX_DIAG_12_RAW,
         };
 
         let stream = res.stream.clone();

deploy/README.md

@@ -37,6 +37,18 @@ deploy on every run. Nothing secret lives here.
   advances and `miner_gpu_devices` matches the matrix.
 - **Binaries must be built on `cuda-13.0`** (Fedora 43, like the hosts). The
   `rust` runner is Fedora 44; its binaries fail with `GLIBC_2.43 not found`.
+- **The GPU profile lives and dies with the unit.** `quantus-miner.service`
+  applies the power limit and 810 MHz memory lock in `ExecStartPre` and resets
+  to driver defaults in `ExecStopPost`, via `nvidia-power-limit.sh` from
+  lair/quantus. The lock cuts LLM inference to a sixth of its speed with nothing
+  reporting as throttled, so it must not outlive the miner. Stopping the unit
+  therefore also resets the clocks — the benchmark re-applies the profile for
+  its window.
+- **An inference host is left alone.** When `neuron.service` is enabled on a
+  host, the deploy ships the miner's files but does not enable or start it, and
+  validate passes without checking hashing. neuron conflicts with the miner, so
+  starting it would stop inference. Switch hosts deliberately:
+  `systemctl disable --now neuron && systemctl enable --now quantus-miner`.
 - **A benchmark (`bench.yaml`) stops the miner on its host.** A deploy landing
   on the same host during a measurement would find the unit down and start it
   mid-window. Both are serialised only by not running them together.

deploy/infra-setup.sh

@@ -110,6 +110,11 @@ gitea_ci ALL=(root) NOPASSWD: /usr/bin/systemctl restart quantus-miner.service
 # mining for a measurement window and resume it afterwards.
 gitea_ci ALL=(root) NOPASSWD: /usr/bin/systemctl stop quantus-miner.service
 gitea_ci ALL=(root) NOPASSWD: /usr/bin/systemctl start quantus-miner.service
+# The benchmark measures under the miner's GPU profile, which stopping the unit
+# resets (its ExecStopPost). bench-on-host.sh re-applies it for the window and
+# restores whichever profile the host was in afterwards.
+gitea_ci ALL=(root) NOPASSWD: /usr/local/bin/nvidia-power-limit.sh miner
+gitea_ci ALL=(root) NOPASSWD: /usr/local/bin/nvidia-power-limit.sh default
 # Rollback: the deploy keeps the previous binary as .prev (cp) and restores it
 # with install when validate fails; cp back would hit "Text file busy" on the
 # running binary, install unlinks the destination first. Deploying from main means a bad commit reaches production;

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

deploy/quantus-miner.service

@@ -9,8 +9,13 @@
 #     finds no adapter, silently falling back to nothing.
 #
 # Poseidon2-over-Goldilocks is compute-bound in registers, not memory-bound, so
-# there is no tuning knob here worth more than the GPU power limit — which is a
-# host-level concern (nvidia-smi -pl), not a unit-file one.
+# the tuning that matters is the GPU profile: a lower power limit and memory
+# locked at 810 MHz (lair/quantus readme, "GPU memory clock"). That profile is
+# applied here, for the lifetime of this unit, and reset when it stops — it must
+# never outlive the miner, because the same lock cuts LLM inference on the host
+# to a sixth of its speed. The script and its per-host config
+# (/etc/default/nvidia-power-limit) are installed by lair/quantus's
+# infra-setup.sh; neuron.service on the same host conflicts with this unit.
 #
 # --cpu-workers 0 is deliberate: a current-gen desktop CPU contributes ~4 MH/s
 # next to a 4090's ~183 MH/s, for a couple hundred watts. Measured, not assumed.
@@ -37,6 +42,12 @@ ExecStart=/usr/local/bin/quantus-miner serve \
     --metrics-port 9900
 EnvironmentFile=/etc/quantus-miner/miner.env
 
+# '+' runs these as root; the miner itself stays unprivileged. ExecStopPost runs
+# however the service ends — stop, crash, failed start — so the lock cannot be
+# left behind by a miner that is no longer running.
+ExecStartPre=+/usr/local/bin/nvidia-power-limit.sh miner
+ExecStopPost=+/usr/local/bin/nvidia-power-limit.sh default
+
 # The node may be down, still syncing, or mid-redeploy; reconnecting is normal
 # operation, not an error condition.
 Restart=always