engine-cuda: per-arch carry path, +20.5% on sm_120

An `asm` block is opaque to nvcc's optimiser, so carry arithmetic written
as inline PTX blocks common-subexpression elimination and strength
reduction across neighbouring field operations. On Blackwell that costs a
quarter of the kernel. On Ada and Ampere the hand-written sequence still
wins, by a similar margin in the other direction, so the default is now
chosen per device pass.

Whole-kernel instructions (ptxas + cuobjdump, CUDA 13.0) and measured
hashrate, three interleaved rounds each:

                    sm_120                sm_86/89
  PTX_CARRY=1   37,432   1127.1 MH/s   37,027   benjy 461.6  quadbrat 85.5
  PTX_CARRY=0   27,837   1358.0 MH/s   41,196   benjy 401.5  quadbrat 76.3

sm_86/89 keep the exact configuration they had, so those cards are
unchanged by construction -- identical macro values, identical 37,027
instructions. Parity 60/60 against the CPU on beast.

LAIR_PTX_ACC is now 1 everywhere rather than following LAIR_PTX_CARRY. In
the Acc accumulators the add-with-carry pair is the whole operation, so
there is nothing around it for the optimiser to fold; on sm_120 with the
carry path off, Acc PTX on measured 1333.3 against 1285.8 with it off.

PR #17 measured the PTX carry path as a win on every card. That was true
of the kernel it was written for, which spilled 104 bytes per thread;
once the spills went the trade reversed on Blackwell only.

Found by profiling rather than reading: ncu puts us at 43,700 executed
instructions per nonce against the rival binary's 22,229 on the same
card, a 1.97x ratio that matches the 2.05x throughput gap, with 82% of
the excess on the ALU pipe rather than the multiply pipe. This lands
35,785. The remaining gap is still ALU: 26,586 against their 14,413,
while the multiply pipe is now close to parity at 8,405 against 7,030.

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

.gitea/workflows/bench.yaml

@@ -0,0 +1,143 @@
+---
+# GPU hashrate and parity measurement on a fleet mining host (quantus/miner#2).
+#
+# Runner containers on the GPU hosts have no device passthrough (gongfoo gives
+# them `devices: None`), so the benchmark cannot run inside the runner. The
+# binary is built on a runner and executed on the host over ssh as gitea_ci,
+# which has the GPU device nodes (0666) and the Vulkan ICD, exactly like the
+# miner. The host's quantus-miner.service is stopped for the window and started
+# again afterwards, under a trap, so a failed run cannot leave the host idle.
+#
+# Benjy (4090, dedicated) is the reference card; beast (2x 5090) also serves
+# inference and is only benchmarked by manual dispatch.
+name: bench
+
+on:
+  workflow_dispatch:
+    inputs:
+      host:
+        description: mining host to measure
+        default: benjy.hanzalova.internal
+        type: choice
+        options:
+          - benjy.hanzalova.internal
+          - quadbrat.hanzalova.internal
+          - beast.hanzalova.internal
+      duration_secs:
+        description: seconds per timed window
+        default: "30"
+      runs:
+        description: timed windows (median reported)
+        default: "5"
+      batch_size:
+        description: nonces per GPU batch
+        default: "1000000"
+      workers:
+        description: worker threads (one per card)
+        default: "1"
+  pull_request:
+    paths:
+      - crates/engine-gpu/**
+      - crates/engine-cuda/**
+      - crates/engine-cpu/**
+      - crates/pow-core/**
+      - crates/miner-service/**
+      - crates/bench-harness/**
+      - Cargo.lock
+      - .gitea/workflows/bench.yaml
+
+env:
+  CARGO_TERM_COLOR: always
+  BENCH_HOST: ${{ github.event.inputs.host || 'benjy.hanzalova.internal' }}
+  DURATION: ${{ github.event.inputs.duration_secs || '30' }}
+  RUNS: ${{ github.event.inputs.runs || '5' }}
+  BATCH: ${{ github.event.inputs.batch_size || '1000000' }}
+  WORKERS: ${{ github.event.inputs.workers || '1' }}
+
+# One measurement per host at a time; a second one would share the card.
+concurrency:
+  group: bench-${{ github.event.inputs.host || 'benjy.hanzalova.internal' }}
+  cancel-in-progress: false
+
+jobs:
+  build:
+    # cuda-13.0, not rust: the rust image is Fedora 44 and the mining hosts are
+    # Fedora 43, so a binary built there needs a glibc the hosts do not have
+    # (observed: "GLIBC_2.43 not found"). cuda-13.0 is Fedora 43 based, matches
+    # the hosts, and is what the deploy in #8 builds on anyway.
+    runs-on: cuda-13.0
+    steps:
+      - uses: actions/checkout@v4
+      - name: build quantus-bench
+        env:
+          MINER_BUILD_SHA: ${{ github.event.pull_request.head.sha || github.sha }}
+        run: cargo build --release --locked -p bench-harness
+      - uses: actions/upload-artifact@v3
+        with:
+          name: quantus-bench
+          path: target/release/quantus-bench
+
+  measure:
+    runs-on: fedora-43
+    needs: build
+    steps:
+      - uses: actions/checkout@v4
+      - uses: actions/download-artifact@v3
+        with: { name: quantus-bench, path: _bin }
+
+      - name: write ssh key
+        run: |
+          set -euo pipefail
+          install -d -m 0700 ~/.ssh
+          printf '%s\n' "${{ secrets.RSYNC_SSH_KEY }}" > ~/.ssh/id_gitea_ci
+          chmod 0600 ~/.ssh/id_gitea_ci
+
+      - name: measure on ${{ env.BENCH_HOST }}
+        run: |
+          set -euo pipefail
+          SSHOPTS="-i $HOME/.ssh/id_gitea_ci -o StrictHostKeyChecking=accept-new"
+          run() { ssh $SSHOPTS gitea_ci@"$BENCH_HOST" "$@"; }
+
+          # Per-run file names: two runs may overlap on the host until the
+          # host-side flock serialises them, and scp over a running binary
+          # fails with ETXTBSY.
+          id="${{ github.run_id }}"
+          dir=/var/lib/gitea_ci/bench
+          run "install -d -m 0750 $dir"
+          scp $SSHOPTS -q _bin/quantus-bench gitea_ci@"$BENCH_HOST":$dir/quantus-bench.$id
+          scp $SSHOPTS -q crates/bench-harness/bench-on-host.sh gitea_ci@"$BENCH_HOST":$dir/bench-on-host.$id.sh
+          run "chmod 0755 $dir/quantus-bench.$id"
+          cleanup() { ssh $SSHOPTS gitea_ci@"$BENCH_HOST" "rm -f $dir/quantus-bench.$id $dir/bench-on-host.$id.sh $dir/record.$id.json"; }
+          trap cleanup EXIT
+
+          label="${{ github.event.pull_request.number && format('pr-{0}', github.event.pull_request.number) || github.ref_name }}"
+          # The host script pauses the miner, holds the per-host lock, measures,
+          # and resumes the miner under its own trap.
+          run "bash $dir/bench-on-host.$id.sh $dir/quantus-bench.$id $DURATION $RUNS $BATCH $WORKERS $label $dir/record.$id.json" | tee bench.log
+          # The summary is everything from the harness's markdown header on.
+          sed -n '/^## quantus-bench/,$p' bench.log > bench.md
+          scp $SSHOPTS -q gitea_ci@"$BENCH_HOST":$dir/record.$id.json record.json
+          { echo; cat bench.md; } >> "$GITHUB_STEP_SUMMARY"
+
+      - uses: actions/upload-artifact@v3
+        if: always()
+        with:
+          name: bench-record-${{ env.BENCH_HOST }}
+          path: |
+            record.json
+            bench.md
+
+      - name: comment on the pull request
+        if: github.event_name == 'pull_request'
+        run: |
+          set -euo pipefail
+          body=$(python3 - <<'PY'
+          import json, pathlib
+          print(json.dumps({"body": pathlib.Path("bench.md").read_text()}))
+          PY
+          )
+          curl -fsS -X POST \
+            -H "Authorization: token ${{ secrets.GITHUB_TOKEN }}" \
+            -H "Content-Type: application/json" \
+            "${GITHUB_API_URL}/repos/${GITHUB_REPOSITORY}/issues/${{ github.event.pull_request.number }}/comments" \
+            -d "$body" > /dev/null

.gitea/workflows/ci.yml

@@ -0,0 +1,81 @@
+---
+# Lint, build and test on the fleet's `rust` runner (architecture/gitea-runners.md).
+#
+# Replaces origin's .github/workflows/ci.yml, which targeted GitHub-hosted
+# runners that do not exist here. Gitea Actions also reads .github/workflows,
+# so origin's files are deleted rather than left inert; on every origin merge
+# `git rm -r .github/workflows` is the standing resolution (quantus/miner#1).
+#
+# Deliberately not here: origin's CPU benchmark job (meaningless on a shared
+# 4-CPU runner; the real gate is the GPU harness, quantus/miner#2) and taplo
+# (not on the runner image; add it to runner-rust per gitea-runners.md §5
+# rather than `cargo install` on every run).
+name: ci
+
+on:
+  pull_request:
+  push:
+    branches: [main, origin-main]
+
+concurrency:
+  group: ci-${{ github.event.pull_request.number || github.ref }}
+  cancel-in-progress: true
+
+env:
+  CARGO_INCREMENTAL: "0"
+  CARGO_TERM_COLOR: always
+
+# rust-toolchain pins 1.93.0. rustup on the runner auto-installs the toolchain
+# on first use but NOT the components the file lists (observed: "cargo-fmt is
+# not installed for the toolchain"), so each job that needs one adds it
+# explicitly; the call is a no-op once present.
+
+jobs:
+  fmt:
+    runs-on: rust
+    steps:
+      - uses: actions/checkout@v4
+      - name: toolchain
+        run: rustup component add rustfmt
+      - name: rustfmt
+        run: cargo fmt --all -- --check
+
+  clippy:
+    runs-on: rust
+    needs: fmt
+    steps:
+      - uses: actions/checkout@v4
+      - name: toolchain
+        run: rustup component add clippy
+      - name: clippy (all targets, all features, warnings denied)
+        # Same invocation as clippy.sh so local and CI agree.
+        run: cargo clippy --workspace --all-targets --all-features --locked -- -D warnings
+
+  test:
+    runs-on: rust
+    needs: fmt
+    steps:
+      - uses: actions/checkout@v4
+      - name: build
+        run: cargo build --locked --workspace
+      - name: test
+        # GPU tests skip themselves when no adapter is present; the runner has none.
+        run: cargo test --locked --workspace
+      - name: version carries the commit
+        # build.rs embeds the SHA; a binary that says "unknown" cannot be
+        # matched to a deploy (quantus/miner#8).
+        run: |
+          set -euo pipefail
+          v=$(./target/debug/quantus-miner --version)
+          echo "$v"
+          case "$v" in
+            *unknown*) echo "build SHA not embedded" >&2; exit 1 ;;
+          esac
+
+  doc:
+    runs-on: rust
+    needs: fmt
+    steps:
+      - uses: actions/checkout@v4
+      - name: doc
+        run: cargo doc --locked --workspace --no-deps --all-features

.gitea/workflows/deploy.yaml

@@ -0,0 +1,382 @@
+---
+# Deploy — or validate — the miner on the fleet's mining hosts (quantus/miner#8).
+#
+# Every push to main builds the binary and lands it on each host in the matrix,
+# validated, with rollback. This workflow is the source of infra truth for the
+# miner: hosts and per-host settings live in the `deploy` matrix and nowhere
+# else. The node it attaches to is deployed by quantus/chain; the fleet's
+# monitoring, GPU power limits and nvidia metrics stay in lair/quantus.
+#
+# Build on cuda-13.0: it is Fedora 43 like the hosts (the `rust` image is
+# Fedora 44 and its binaries fail on the hosts with GLIBC_2.43 not found), and
+# it is the only runner with nvcc for the CUDA engine (#3). Deploy on
+# fedora-43: ssh + rsync are on every runner (gitea-runners.md §3).
+name: deploy
+
+on:
+  push:
+    branches: [main]
+    paths-ignore:
+      - "**.md"
+      - .gitea/workflows/ci.yml
+      - .gitea/workflows/bench.yaml
+  workflow_dispatch:
+    inputs:
+      mode:
+        description: "deploy (apply) or validate (check only, no changes)"
+        required: false
+        default: deploy
+        type: choice
+        options: [deploy, validate]
+      chain:
+        description: "chain id under the node's chains/ directory to take miner credentials from (overrides CHAIN); the node deploy's validate step prints it"
+        required: false
+
+# Never half-apply two deploys at once. (Not relied on for correctness of a
+# single host: the restart is checksum-gated and the binary is rsynced atomically.)
+concurrency:
+  group: deploy-miner
+  cancel-in-progress: false
+
+env:
+  # The chain spec: the miner credential path on the node host derives from it.
+  # Must agree with lair/quantus's node deploy. Mainnet since 2026-09-09; the
+  # node writes the credentials under chains/mainnet/.
+  CHAIN: mainnet
+  MINER_LINK_PORT: "9833"
+  MINER_METRICS_PORT: "9900"
+  # Fleet Prometheus host; the miner's exporter is unauthenticated and bound to
+  # 0.0.0.0, so this is the only host allowed to reach it.
+  METRICS_HOST: golgafrinchans.kosherinata.internal
+  CARGO_TERM_COLOR: always
+
+jobs:
+  build:
+    runs-on: cuda-13.0
+    steps:
+      - uses: actions/checkout@v4
+      - name: build quantus-miner
+        env:
+          # Embedded by crates/miner-cli/build.rs into --version; validate below
+          # asserts the host runs exactly this commit.
+          MINER_BUILD_SHA: ${{ github.sha }}
+        run: |
+          set -euo pipefail
+          cargo build --release --locked -p miner-cli
+          ./target/release/quantus-miner --version
+      - uses: actions/upload-artifact@v3
+        with:
+          name: quantus-miner
+          path: target/release/quantus-miner
+
+  deploy:
+    runs-on: fedora-43
+    needs: build
+    strategy:
+      fail-fast: false        # one host's failure must not abort the other
+      matrix:
+        include:
+          # `node` is the host whose QUIC control channel this miner attaches
+          # to, and whose reward address therefore receives what it earns. The
+          # node's deploy (quantus/chain) must list this host in its `miners`
+          # 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.
+          # `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.
+          - host: benjy.hanzalova.internal
+            node: bob.hanzalova.internal
+            gpu_devices: "1"        # 1x RTX 4090 (sm_89); reference card for #2
+            kernel: cuda
+          - host: quadbrat.hanzalova.internal
+            node: bob.hanzalova.internal
+            gpu_devices: "1"        # 1x RTX 3060 (sm_86)
+            kernel: cuda
+          - host: beast.hanzalova.internal
+            node: bob.hanzalova.internal
+            gpu_devices: "2"        # 2x RTX 5090 (sm_120)
+            kernel: cuda
+    steps:
+      - uses: actions/checkout@v4
+      - uses: actions/download-artifact@v3
+        with: { name: quantus-miner, path: _bin }
+
+      - name: write ssh key
+        run: |
+          set -euo pipefail
+          install -d -m 0700 ~/.ssh
+          printf '%s\n' "${{ secrets.RSYNC_SSH_KEY }}" > ~/.ssh/id_gitea_ci
+          chmod 0600 ~/.ssh/id_gitea_ci
+
+      - name: reachability
+        run: |
+          set -euo pipefail
+          ssh -i ~/.ssh/id_gitea_ci -o StrictHostKeyChecking=accept-new \
+            gitea_ci@${{ matrix.host }} hostname -f
+
+      - name: preflight — sudoers covers this deploy
+        run: |
+          set -euo pipefail
+          SSHOPTS="-i $HOME/.ssh/id_gitea_ci -o StrictHostKeyChecking=accept-new"
+          # Every path deploy/infra-setup.sh grants must already be permitted on
+          # the host. Compare up front, from the script itself so the two cannot
+          # drift, and name the missing paths instead of failing forty lines into
+          # an rsync with "sudo: a password is required".
+          sed -n "/quantus-miner_gitea_ci.tmp/,/^SUDO$/p" deploy/infra-setup.sh \
+            | grep '^gitea_ci ALL=' | grep -oE '(/etc|/usr|/var)/[^ ]*' | sort -u > expected-paths.txt
+          ssh $SSHOPTS gitea_ci@${{ matrix.host }} 'sudo -n -l' \
+            | grep -oE '(/etc|/usr|/var)/[^ ]*' | sort -u > permitted-paths.txt
+          comm -23 expected-paths.txt permitted-paths.txt > missing-paths.txt
+          if [ -s missing-paths.txt ]; then
+            echo "the miner sudoers on ${{ matrix.host }} is out of date." >&2
+            echo "not permitted, but this deploy needs them:" >&2
+            sed 's/^/  /' missing-paths.txt >&2
+            echo "" >&2
+            echo "run: ./deploy/infra-setup.sh --pubkey ~/.ssh/id_gitea_ci.pub" >&2
+            exit 1
+          fi
+          echo "sudoers covers all $(wc -l < expected-paths.txt) paths this deploy needs"
+
+      - name: deploy miner
+        id: deploy
+        if: ${{ github.event.inputs.mode != 'validate' }}
+        env:
+          GPU_DEVICES: ${{ matrix.gpu_devices }}
+        run: |
+          set -euo pipefail
+          SSHOPTS="-i $HOME/.ssh/id_gitea_ci -o StrictHostKeyChecking=accept-new"
+          nrun() { ssh $SSHOPTS gitea_ci@"${{ matrix.node }}" "$@"; }
+          run()  { ssh $SSHOPTS gitea_ci@"${{ matrix.host }}" "$@"; }
+
+          # -c (checksum), not rsync's default size+mtime quick check: the
+          # artifact is freshly built every run so mtimes always differ, and the
+          # default heuristic would report a change on every deploy. Restart is
+          # gated on an itemised content difference only.
+          RESTART=0
+          push() {
+            local out
+            out=$(rsync -e "ssh $SSHOPTS" --rsync-path='sudo rsync' -ic "$@")
+            if [ -n "$out" ]; then
+              RESTART=1
+              printf '%s\n' "$out" | sed 's/^/    changed: /'
+            fi
+          }
+
+          # 1. service account + dirs
+          push --mkpath --chmod=F0644 \
+            deploy/quantus-miner.sysusers.conf \
+            gitea_ci@"${{ matrix.host }}":/etc/sysusers.d/quantus-miner.conf
+          run sudo systemd-sysusers
+          run sudo install -d -o root -g quantus-miner -m 0750 /etc/quantus-miner
+          run sudo install -d -o quantus-miner -g quantus-miner -m 0750 /var/lib/quantus-miner
+
+          # 2. The miner's credentials are GENERATED BY THE NODE on first start
+          #    and regenerate if the node's base-path is ever wiped. Copying them
+          #    on every deploy is what makes that self-healing instead of a
+          #    silent auth failure. They pass through the runner in memory,
+          #    never the workspace.
+          umask 077
+          tmp=$(mktemp -d)
+          trap 'rm -rf "$tmp"' EXIT
+          nrun sudo cat /var/lib/quantus-node/chains/${{ github.event.inputs.chain || env.CHAIN }}/miner-auth-token \
+            > "$tmp/miner-auth-token"
+          nrun sudo cat /var/lib/quantus-node/chains/${{ github.event.inputs.chain || env.CHAIN }}/miner-tls-cert-sha256 \
+            > "$tmp/miner-tls-cert-sha256"
+          test -s "$tmp/miner-auth-token"      || { echo "node auth token empty — has ${{ matrix.node }} started?" >&2; exit 1; }
+          test -s "$tmp/miner-tls-cert-sha256" || { echo "node TLS pin empty — has ${{ matrix.node }} started?" >&2; exit 1; }
+          push --chown=root:quantus-miner --chmod=F0640 \
+            "$tmp/miner-auth-token" gitea_ci@"${{ matrix.host }}":/etc/quantus-miner/miner-auth-token
+          push --chown=root:quantus-miner --chmod=F0640 \
+            "$tmp/miner-tls-cert-sha256" gitea_ci@"${{ matrix.host }}":/etc/quantus-miner/miner-tls-cert-sha256
+
+          # 3. non-secret runtime config.
+          #    --node-addr parses as a Rust SocketAddr: an IP and a port, with NO
+          #    DNS resolution. Resolve on the MINER host — it is the one
+          #    dialling — and keep the 10.x literal out of the repo.
+          node_addrs=$(run "getent ahostsv4 ${{ matrix.node }}")
+          node_ip=$(awk '{print $1; exit}' <<<"$node_addrs")
+          case "$node_ip" in
+            10.*) echo "node address: ${{ matrix.node }} -> ${node_ip}" ;;
+            *) echo "refusing to point the miner at non-mesh address '${node_ip}'" >&2; exit 1 ;;
+          esac
+          export NODE_ADDR="${node_ip}:${{ env.MINER_LINK_PORT }}"
+
+          python3 - <<'PY'
+          import os, pathlib
+          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"])
+          pathlib.Path("miner.env").write_text(t)
+          PY
+          push --chown=root:quantus-miner --chmod=F0640 \
+            miner.env gitea_ci@"${{ matrix.host }}":/etc/quantus-miner/miner.env
+
+          # 4. binary + unit. Keep the running binary as .prev first so a failed
+          #    validate can put it back (rollback step below).
+          if run test -x /usr/local/bin/quantus-miner; then
+            run sudo cp -p /usr/local/bin/quantus-miner /usr/local/bin/quantus-miner.prev
+            echo "previous binary kept: $(run /usr/local/bin/quantus-miner.prev --version)"
+          fi
+          push --chmod=F0755 _bin/quantus-miner gitea_ci@"${{ matrix.host }}":/usr/local/bin/quantus-miner
+          push --chmod=F0644 deploy/quantus-miner.service \
+            gitea_ci@"${{ matrix.host }}":/etc/systemd/system/quantus-miner.service
+
+          run sudo restorecon -R /usr/local/bin/quantus-miner /etc/quantus-miner /var/lib/quantus-miner
+
+          # 5. firewalld for the exporter, scoped to the scrape host. The miner
+          #    binds metrics on 0.0.0.0 unconditionally.
+          rsync -e "ssh $SSHOPTS" --rsync-path='sudo rsync' -ic --mkpath --chmod=F0644 \
+            deploy/quantus-miner-metrics.xml \
+            gitea_ci@"${{ matrix.host }}":/etc/firewalld/services/quantus-miner-metrics.xml \
+            | sed 's/^/    changed: /'
+          run sudo firewall-cmd --reload
+          zone=$(run sudo firewall-cmd --get-default-zone)
+          metrics_addrs=$(run "getent ahostsv4 ${{ env.METRICS_HOST }}")
+          metrics_ip=$(awk '{print $1; exit}' <<<"$metrics_addrs")
+          case "$metrics_ip" in
+            10.*) echo "scrape source: ${metrics_ip}" ;;
+            *) echo "refusing to expose metrics to non-mesh address '${metrics_ip}'" >&2; exit 1 ;;
+          esac
+          mrich="rule family=ipv4 source address=${metrics_ip}/32 service name=quantus-miner-metrics accept"
+          if run "sudo firewall-cmd --zone=$zone --query-rich-rule='$mrich'"; then
+            echo "firewalld: metrics rich rule already present in ${zone}"
+          else
+            run "sudo firewall-cmd --permanent --zone=$zone --add-rich-rule='$mrich'"
+            run "sudo firewall-cmd --zone=$zone --add-rich-rule='$mrich'"
+          fi
+
+          run sudo systemctl enable quantus-miner.service   # idempotent
+          if [ "$RESTART" = 1 ]; then
+            echo "changes applied — restarting"
+            run sudo systemctl daemon-reload
+            run sudo systemctl restart quantus-miner.service
+          elif run systemctl is-active --quiet quantus-miner.service; then
+            echo "nothing changed and the miner is running — left alone"
+          else
+            echo "nothing changed but the miner is down — starting it"
+            run sudo systemctl restart quantus-miner.service
+          fi
+          echo "restarted=$RESTART" >> "$GITHUB_OUTPUT"
+
+      - name: validate miner
+        id: validate
+        run: |
+          set -euo pipefail
+          SSHOPTS="-i $HOME/.ssh/id_gitea_ci -o StrictHostKeyChecking=accept-new"
+          run() { ssh $SSHOPTS gitea_ci@"${{ matrix.host }}" "$@"; }
+          fail=0
+
+          echo "--- unit (${{ matrix.host }} -> ${{ matrix.node }}) ---"
+          run systemctl is-active quantus-miner.service
+
+          echo "--- version ---"
+          got=$(run /usr/local/bin/quantus-miner --version)
+          echo "installed: ${got}"
+          case "$got" in
+            *"${{ github.sha }}"*) echo "binary is this commit" ;;
+            *)
+              if [ "${{ github.event.inputs.mode }}" = validate ]; then
+                echo "note: installed binary is not this commit (validate mode; not a failure)"
+              else
+                echo "installed binary does NOT carry commit ${{ github.sha }}" >&2; fail=1
+              fi ;;
+          esac
+
+          echo "--- gpu ---"
+          # An `active` miner that found no adapter still looks healthy to
+          # systemd; assert the GPU is enumerated and that the miner sees the
+          # number of devices the matrix says it should.
+          run "nvidia-smi --query-gpu=name,power.draw,utilization.gpu --format=csv,noheader"
+
+          echo "--- hashing ---"
+          # The counter is the only honest evidence this process is doing work
+          # rather than idling on a failed connection. After a restart the miner
+          # exports miner_gpu_devices only once it has connected to the node and
+          # miner_hashes_total only once it has hashed, so first wait for the
+          # gauge to appear (readiness), then require the counter to advance.
+          # Here-strings, not pipes: a pipe whose reader exits early SIGPIPEs the
+          # writer and pipefail turns that into exit 141.
+          scrape() { ssh $SSHOPTS gitea_ci@"${{ matrix.host }}" "curl -fsS http://127.0.0.1:${{ env.MINER_METRICS_PORT }}/metrics" || true; }
+          deadline=$((SECONDS + 90))
+          devs=""
+          while [ $SECONDS -lt $deadline ]; do
+            m=$(scrape)
+            devs=$(awk '/^miner_gpu_devices /{print $2; exit}' <<<"$m")
+            [ -n "$devs" ] && break
+            sleep 5
+          done
+          if [ -z "$devs" ]; then
+            echo "  miner did not connect to ${{ matrix.node }} within 90s (no miner_gpu_devices exported)" >&2; fail=1
+          elif [ "${devs%.*}" != "${{ matrix.gpu_devices }}" ]; then
+            echo "  miner_gpu_devices is ${devs}, matrix says ${{ matrix.gpu_devices }}" >&2; fail=1
+          else
+            echo "  ok   connected; miner_gpu_devices ${devs%.*}"
+          fi
+          # The build-info gauge is what ties every other metric to a commit
+          # (quantus/miner#9). In deploy mode it must carry this commit.
+          if grep -q "^miner_build_info{.*commit=\"${{ github.sha }}\"" <<<"$m"; then
+            echo "  ok   miner_build_info carries ${{ github.sha }}"
+          elif [ "${{ github.event.inputs.mode }}" != validate ]; then
+            echo "  miner_build_info does not carry commit ${{ github.sha }}" >&2; fail=1
+          fi
+          # The kernel actually running. A wgpu fallback on a host that should
+          # run CUDA passes every other check at a fraction of the hashrate.
+          if grep -q "^miner_device_hashes_total{.*kernel=\"${{ matrix.kernel }}\"" <<<"$m"; then
+            echo "  ok   kernel ${{ matrix.kernel }} on device 0"
+          else
+            echo "  expected kernel ${{ matrix.kernel }}, exported series:" >&2
+            grep "^miner_device_hashes_total" <<<"$m" >&2 || echo "  (none yet)" >&2
+            if [ "${{ github.event.inputs.mode }}" != validate ]; then fail=1; fi
+          fi
+
+          deadline=$((SECONDS + 120))
+          h1=""; h2=""; ok=0
+          while [ $SECONDS -lt $deadline ]; do
+            m=$(scrape)
+            h=$(awk '/^miner_hashes_total /{print $2; exit}' <<<"$m")
+            if [ -n "$h" ]; then
+              if [ -z "$h1" ]; then
+                h1="$h"
+              elif [ "${h%.*}" -gt "${h1%.*}" ]; then
+                h2="$h"; ok=1; break
+              fi
+            fi
+            sleep 10
+          done
+          if [ "$ok" = 1 ]; then
+            echo "  ok   miner_hashes_total ${h1%.*} -> ${h2%.*}"
+          else
+            echo "  miner_hashes_total did not advance within 120s (${h1:-none} -> ${h2:-none})" >&2
+            fail=1
+          fi
+
+          exit $fail
+
+      - name: rollback
+        # Only when this run replaced the binary and validate then failed.
+        if: ${{ failure() && steps.deploy.outputs.restarted == '1' }}
+        run: |
+          set -euo pipefail
+          SSHOPTS="-i $HOME/.ssh/id_gitea_ci -o StrictHostKeyChecking=accept-new"
+          run() { ssh $SSHOPTS gitea_ci@"${{ matrix.host }}" "$@"; }
+          if run test -x /usr/local/bin/quantus-miner.prev; then
+            echo "validate failed after a restart — restoring the previous binary"
+            # install, not cp: cp writes in place and fails with "Text file busy"
+            # on a running binary; install unlinks the destination first (the
+            # same reason rsync's temp-file-and-rename works for the push).
+            run sudo install -m 0755 /usr/local/bin/quantus-miner.prev /usr/local/bin/quantus-miner
+            run sudo restorecon -R /usr/local/bin/quantus-miner /etc/quantus-miner /var/lib/quantus-miner
+            run sudo systemctl restart quantus-miner.service
+            echo "restored: $(run /usr/local/bin/quantus-miner --version)"
+          else
+            echo "no previous binary to restore" >&2
+          fi
+
+      - name: journal
+        if: always()
+        run: |
+          ssh -i ~/.ssh/id_gitea_ci -o StrictHostKeyChecking=accept-new \
+            gitea_ci@${{ matrix.host }} journalctl -u quantus-miner.service -n 60 --no-pager

.github/actions/disk/action.yml

@@ -1,15 +0,0 @@
----
-name: free disk space
-description: when rust compiling, free up some disk space
-
-runs:
-    using: composite
-    steps:
-        - name: free disk space
-          uses: jlumbroso/free-disk-space@54081f138730dfa15788a46383842cd2f914a1be # 1.3.1
-          with:
-              android: true
-              dotnet: false
-              haskell: false
-              large-packages: false
-              swap-storage: false

.github/workflows/ci.yml

@@ -1,96 +0,0 @@
----
-name: CI
-
-on:
-  pull_request:
-    paths-ignore:
-      - "docs/**"
-      - "*.md"
-      - "LICENSE"
-  push:
-    branches:
-      - main
-    paths-ignore:
-      - "docs/**"
-      - "*.md"
-      - "LICENSE"
-
-concurrency:
-  group: ${{ github.workflow }}-${{ github.event.pull_request.number || github.ref }}
-  cancel-in-progress: true
-
-env:
-  CARGO_INCREMENTAL: 0
-  CARGO_TERM_COLOR: always
-
-jobs:
-  fast-checks:
-    name: 🏁 Fast Checks (Format)
-    runs-on: ubuntu-latest
-    steps:
-      - uses: actions/checkout@v4
-      - name: setup rust
-        uses: actions-rust-lang/setup-rust-toolchain@v1
-        with:
-          components: rustfmt
-      - name: install taplo
-        run: cargo install taplo-cli --locked
-      - name: Run format checks
-        run: |
-          taplo format --check --config taplo.toml
-          cargo fmt --all -- --check
-        timeout-minutes: 5
-
-  build-and-test:
-    name: 🛠️ Build & Test
-    needs: fast-checks
-    runs-on: ubuntu-latest
-    steps:
-      - uses: actions/checkout@v4
-      - name: setup rust
-        uses: actions-rust-lang/setup-rust-toolchain@v1
-      - name: compile (gpu)
-        run: |
-          cargo build --locked --workspace
-        timeout-minutes: 90
-      - name: test (gpu)
-        run: |
-          cargo test --locked --workspace
-        timeout-minutes: 15
-
-  analysis:
-    name: 🤖 Analysis (Clippy & Doc)
-    needs: fast-checks
-    runs-on: ubuntu-latest
-    steps:
-      - uses: actions/checkout@v4
-      - name: setup rust
-        uses: actions-rust-lang/setup-rust-toolchain@v1
-        with:
-          components: clippy
-      - name: clippy (all features)
-        run: cargo clippy --locked --workspace --all-features
-        timeout-minutes: 30
-      - name: doc
-        run: cargo doc --locked --workspace --no-deps --all-features
-        timeout-minutes: 15
-
-  benchmark:
-    name: 🏃 Benchmark
-    needs: fast-checks
-    runs-on: ubuntu-latest
-    steps:
-      - uses: actions/checkout@v4
-      - name: setup rust
-        uses: actions-rust-lang/setup-rust-toolchain@v1
-      - name: build benchmark binary
-        run: cargo build -p miner-cli --release
-        timeout-minutes: 60
-      - name: upload linux release binary
-        uses: actions/upload-artifact@v4
-        with:
-          name: quantus-miner-linux-x86_64
-          path: target/release/quantus-miner
-      - name: run cpu benchmark
-        run: ./target/release/quantus-miner benchmark --cpu-workers 2 --duration 5
-        timeout-minutes: 10

.github/workflows/dependency-cooldown-audit.yml

@@ -1,13 +0,0 @@
-name: Dependency cooldown audit
-
-on:
-  schedule:
-    - cron: '0 6 * * 1'
-  workflow_dispatch:
-
-permissions:
-  contents: read
-
-jobs:
-  audit:
-    uses: Quantus-Network/shared-workflows/.github/workflows/dependency-cooldown-audit.yml@v1

.github/workflows/dependency-cooldown.yml

@@ -1,14 +0,0 @@
-name: Dependency cooldown
-
-on:
-  pull_request:
-    # labeled/unlabeled/edited are required so the gate re-runs when an
-    # emergency bypass label or reason is added.
-    types: [opened, synchronize, reopened, labeled, unlabeled, edited]
-
-permissions:
-  contents: read
-
-jobs:
-  dependency-cooldown:
-    uses: Quantus-Network/shared-workflows/.github/workflows/dependency-cooldown.yml@v1

.github/workflows/release-proposal.yml

@@ -1,199 +0,0 @@
----
-name: Release Proposal
-
-env:
-  CARGO_TERM_COLOR: always
-
-on:
-  workflow_dispatch:
-    inputs:
-      target_branch:
-        description: "Target branch for the PR (default: main)"
-        required: false
-        type: string
-        default: "main"
-      version_type:
-        description: "Type of version bump"
-        required: true
-        default: "patch"
-        type: choice
-        options:
-          - patch
-          - minor
-          - major
-          - custom
-      custom_version:
-        description: 'Custom version string (e.g., 1.2.3). Only used if version_type is "custom".'
-        required: false
-      is_draft:
-        description: "Is this a draft release?"
-        required: true
-        type: boolean
-        default: false
-
-jobs:
-  calculate-next-version:
-    name: 🧮 Calculate Next Version
-    runs-on: ubuntu-latest
-    outputs:
-      new_version: ${{ steps.versioner.outputs.new_version }}
-      new_tag: ${{ steps.versioner.outputs.new_tag }}
-      commit_sha_short: ${{ steps.vars.outputs.commit_sha_short }}
-      source_branch: ${{ steps.vars.outputs.source_branch }}
-    steps:
-      - name: Checkout code
-        uses: actions/checkout@v4
-        with:
-          fetch-depth: 0
-          fetch-tags: true
-
-      - name: Get current branch and commit SHA
-        id: vars
-        run: |
-          echo "commit_sha_short=$(git rev-parse --short HEAD)" >> $GITHUB_OUTPUT
-          echo "source_branch=$(git rev-parse --abbrev-ref HEAD)" >> $GITHUB_OUTPUT
-
-      - name: Get latest tag
-        id: latest_tag
-        run: |
-          # Get all version tags and sort them by version
-          latest_semver_tag=$(git tag -l "v[0-9]*.[0-9]*.[0-9]*" | sort -V | tail -n 1)
-
-          # If no tags found, use default
-          if [ -z "$latest_semver_tag" ]; then
-            latest_semver_tag="v0.0.0"
-          fi
-
-          echo "latest_tag_found=$latest_semver_tag" >> $GITHUB_OUTPUT
-          echo "Latest semantic version tag found: $latest_semver_tag"
-
-      - name: Calculate new version
-        id: versioner
-        env:
-          LATEST_TAG: ${{ steps.latest_tag.outputs.latest_tag_found }}
-          VERSION_TYPE: ${{ github.event.inputs.version_type }}
-          CUSTOM_VERSION: ${{ github.event.inputs.custom_version }}
-        run: |
-          # Remove 'v' prefix and any suffix for processing
-          current_version=${LATEST_TAG#v}
-          # Remove any suffix after the version number
-          current_version=$(echo "$current_version" | sed -E 's/-[^-]+$//')
-
-          if [[ "$VERSION_TYPE" == "custom" ]]; then
-            if [[ -z "$CUSTOM_VERSION" ]]; then
-              echo "Error: Custom version is selected but no custom_version string provided."
-              exit 1
-            fi
-            if [[ ! "$CUSTOM_VERSION" =~ ^v ]]; then
-              echo "Error: Custom version string MUST start with 'v' (e.g., v1.2.3)."
-              exit 1
-            fi
-            new_version="$CUSTOM_VERSION"
-          else
-            # Split version and pre-release part
-            IFS='-' read -r version_core prerelease_part <<< "$current_version"
-            IFS='.' read -r major minor patch <<< "$version_core"
-
-            # Increment based on type
-            if [[ "$VERSION_TYPE" == "major" ]]; then
-              major=$((major + 1))
-              minor=0
-              patch=0
-            elif [[ "$VERSION_TYPE" == "minor" ]]; then
-              minor=$((minor + 1))
-              patch=0
-            elif [[ "$VERSION_TYPE" == "patch" ]]; then
-              patch=$((patch + 1))
-            else
-              echo "Error: Invalid version_type: $VERSION_TYPE"
-              exit 1
-            fi
-            new_version="v$major.$minor.$patch"
-          fi
-
-          echo "New version: $new_version"
-          echo "new_version=$new_version" >> $GITHUB_OUTPUT
-          echo "new_tag=$new_version" >> $GITHUB_OUTPUT
-
-  update-cargo-toml:
-    name: 📝 Update version files
-    needs: calculate-next-version
-    runs-on: ubuntu-latest
-    permissions:
-      contents: write
-      pull-requests: write
-    steps:
-      - name: Checkout code
-        uses: actions/checkout@v4
-        with:
-          fetch-depth: 0
-
-      - name: Setup rust
-        uses: actions-rust-lang/setup-rust-toolchain@v1
-
-      - name: Create version bump branch and PR
-        env:
-          NEW_VERSION: ${{ needs.calculate-next-version.outputs.new_version }}
-          NEW_TAG: ${{ needs.calculate-next-version.outputs.new_tag }}
-          GITHUB_TOKEN: ${{ secrets.ADMIN_PAT }}
-          SOURCE_BRANCH: ${{ needs.calculate-next-version.outputs.source_branch }}
-          TARGET_BRANCH: ${{ github.event.inputs.target_branch }}
-        run: |
-          set -ex
-          new_cargo_version=${NEW_VERSION#v}
-          branch_name="release/${NEW_VERSION}"
-
-          # Create new branch from source branch
-          git checkout "$SOURCE_BRANCH"
-          git checkout -b "$branch_name"
-
-          # Update version in workspace Cargo.toml (safer than cargo set-version)
-          echo "Updating workspace Cargo.toml to version: $new_cargo_version"
-          sed -i -E "s/^version\s*=\s*\"[0-9a-zA-Z.-]+\"/version = \"$new_cargo_version\"/" Cargo.toml
-
-          # Regenerate Cargo.lock with precise updates for our packages only
-          cargo update -p miner-cli --precise "$new_cargo_version"
-          cargo update -p miner-service --precise "$new_cargo_version"
-          cargo update -p pow-core --precise "$new_cargo_version"
-          cargo update -p metrics --precise "$new_cargo_version"
-          cargo update -p miner-telemetry --precise "$new_cargo_version"
-          cargo update -p engine-cpu --precise "$new_cargo_version"
-          cargo update -p engine-gpu --precise "$new_cargo_version"
-
-          # Verify everything compiles correctly
-          cargo check --workspace
-
-          # Commit changes
-          git config user.name "${{ github.actor }}"
-          git config user.email "${{ github.actor }}@users.noreply.github.com"
-
-          git add Cargo.toml Cargo.lock
-          git commit -m "bump version to $NEW_VERSION"
-          git push origin "$branch_name"
-
-          # Prepare PR title and body
-          PR_TITLE="Release $NEW_VERSION"
-          PR_BODY="Automated version bump for release $NEW_VERSION.
-
-          ## Overview
-          - Version bump: ${{ github.event.inputs.version_type }}
-          - Type: ${{ github.event.inputs.version_type }}
-          - Draft: ${{ github.event.inputs.is_draft }}
-
-          ## What changed
-          - Updated version in Cargo.toml and Cargo.lock
-
-          Triggered by workflow run: ${{ github.server_url }}/${{ github.repository }}/actions/runs/${{ github.run_id }}"
-
-          # Prepare labels
-          PR_LABELS="release-proposal"
-          if [[ "${{ github.event.inputs.is_draft }}" == "true" ]]; then
-            PR_LABELS="$PR_LABELS,draft-release"
-          fi
-
-          gh pr create \
-            --title "$PR_TITLE" \
-            --body "$PR_BODY" \
-            --base "$TARGET_BRANCH" \
-            --head "$branch_name" \
-            --label "$PR_LABELS"

.github/workflows/release-publish.yml

@@ -1,157 +0,0 @@
----
-name: Release Publish
-
-on:
-  pull_request:
-    types: [closed]
-    branches:
-      - main
-
-permissions:
-  contents: write
-  packages: write
-
-env:
-  CARGO_TERM_COLOR: always
-
-jobs:
-  create-tag:
-    name: Create Tag
-    if: github.event.pull_request.merged == true && contains(github.event.pull_request.labels.*.name, 'release-proposal')
-    runs-on: ubuntu-latest
-    outputs:
-      version: ${{ steps.extract_version.outputs.version }}
-      tag: ${{ steps.extract_version.outputs.tag }}
-      is_draft: ${{ steps.extract_version.outputs.is_draft }}
-    steps:
-      - name: Checkout code
-        uses: actions/checkout@v4
-        with:
-          fetch-depth: 0
-
-      - name: Extract version from PR title
-        id: extract_version
-        run: |
-          # Extract version from PR title (format: "Release vX.Y.Z")
-          VERSION_TAG=$(echo "${{ github.event.pull_request.title }}" | grep -o 'v[0-9]\+\.[0-9]\+\.[0-9]\+')
-          if [ -z "$VERSION_TAG" ]; then
-            echo "Error: Could not extract version from PR title: ${{ github.event.pull_request.title }}"
-            exit 1
-          fi
-
-          VERSION=${VERSION_TAG#v}
-          echo "version=$VERSION" >> $GITHUB_OUTPUT
-          echo "tag=$VERSION_TAG" >> $GITHUB_OUTPUT
-
-          # Check if this is a draft release
-          if [[ "${{ contains(github.event.pull_request.labels.*.name, 'draft-release') }}" == "true" ]]; then
-            echo "is_draft=true" >> $GITHUB_OUTPUT
-          else
-            echo "is_draft=false" >> $GITHUB_OUTPUT
-          fi
-
-          echo "Extracted version: $VERSION"
-          echo "Extracted tag: $VERSION_TAG"
-
-      - name: Create and push tag
-        run: |
-          git config user.name "${{ github.actor }}"
-          git config user.email "${{ github.actor }}@users.noreply.github.com"
-          git tag -a "${{ steps.extract_version.outputs.tag }}" -m "Release ${{ steps.extract_version.outputs.tag }}"
-          git push origin "${{ steps.extract_version.outputs.tag }}"
-
-  build:
-    needs: create-tag
-    runs-on: ${{ matrix.os }}
-    strategy:
-      matrix:
-        include:
-          # Linux builds
-          - os: ubuntu-latest
-            target: x86_64-unknown-linux-gnu
-            binary_name: quantus-miner
-            asset_name: quantus-miner-linux-x86_64
-          # Windows builds
-          - os: windows-latest
-            target: x86_64-pc-windows-msvc
-            binary_name: quantus-miner.exe
-            asset_name: quantus-miner-windows-x86_64.exe
-          # macOS builds (Intel)
-          - os: macos-15-intel
-            target: x86_64-apple-darwin
-            binary_name: quantus-miner
-            asset_name: quantus-miner-macos-x86_64
-          # macOS builds (Apple Silicon)
-          - os: macos-latest
-            target: aarch64-apple-darwin
-            binary_name: quantus-miner
-            asset_name: quantus-miner-macos-aarch64
-
-    steps:
-      - uses: actions/checkout@v4
-        with:
-          ref: ${{ needs.create-tag.outputs.tag }}
-
-      - name: setup rust
-        uses: actions-rust-lang/setup-rust-toolchain@v1
-        with:
-          target: ${{ matrix.target }}
-
-      - name: build binary
-        shell: bash
-        run: |
-          if [[ -n "${{ matrix.features }}" ]]; then
-            cargo build --release --locked --target ${{ matrix.target }} --features ${{ matrix.features }}
-          else
-            cargo build --release --locked --target ${{ matrix.target }}
-          fi
-
-      - name: prepare binary
-        shell: bash
-        run: |
-          cd target/${{ matrix.target }}/release
-          if [[ "${{ matrix.os }}" == "windows-latest" ]]; then
-            cp ${{ matrix.binary_name }} ${{ matrix.asset_name }}
-          else
-            cp ${{ matrix.binary_name }} ${{ matrix.asset_name }}
-            strip ${{ matrix.asset_name }}
-          fi
-
-      - name: upload binary artifact
-        uses: actions/upload-artifact@v4
-        with:
-          name: ${{ matrix.asset_name }}
-          path: target/${{ matrix.target }}/release/${{ matrix.asset_name }}
-
-  release:
-    needs: [create-tag, build]
-    runs-on: ubuntu-latest
-    steps:
-      - uses: actions/checkout@v4
-        with:
-          ref: ${{ needs.create-tag.outputs.tag }}
-
-      - name: download all artifacts
-        uses: actions/download-artifact@v4
-        with:
-          path: artifacts
-
-      - name: create release
-        env:
-          GITHUB_TOKEN: ${{ secrets.GITHUB_TOKEN }}
-          TAG: ${{ needs.create-tag.outputs.tag }}
-          IS_DRAFT: ${{ needs.create-tag.outputs.is_draft }}
-        run: |
-          # Prepare draft flag
-          if [[ "$IS_DRAFT" == "true" ]]; then
-            DRAFT_FLAG="--draft"
-          else
-            DRAFT_FLAG=""
-          fi
-
-          # Create release with all artifacts
-          gh release create "$TAG" \
-            --title "Release $TAG" \
-            --generate-notes \
-            $DRAFT_FLAG \
-            artifacts/*/quantus-miner-*

Cargo.lock

@@ -107,7 +107,7 @@ version = "0.38.0+1.3.281"
 source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "0bb44936d800fea8f016d7f2311c6a4f97aebd5dc86f09906139ec848cf3a46f"
 dependencies = [
- "libloading",
+ "libloading 0.8.9",
 ]
 
 [[package]]
@@ -122,6 +122,23 @@ version = "0.21.7"
 source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "9d297deb1925b89f2ccc13d7635fa0714f12c87adce1c75356b39ca9b7178567"
 
+[[package]]
+name = "bench-harness"
+version = "4.0.2"
+dependencies = [
+ "clap",
+ "engine-cpu",
+ "engine-cuda",
+ "engine-gpu",
+ "env_logger",
+ "log",
+ "pow-core",
+ "primitive-types 0.13.1",
+ "rand 0.9.2",
+ "serde",
+ "serde_json",
+]
+
 [[package]]
 name = "bit-set"
 version = "0.8.0"
@@ -524,6 +541,15 @@ dependencies = [
  "typenum",
 ]
 
+[[package]]
+name = "cudarc"
+version = "0.19.9"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "804764d10e844da09765a7b2ca9641a0851523d1702efb0d7299d73e31b86e80"
+dependencies = [
+ "libloading 0.9.0",
+]
+
 [[package]]
 name = "data-encoding"
 version = "2.10.0"
@@ -586,6 +612,19 @@ dependencies = [
  "rand 0.9.2",
 ]
 
+[[package]]
+name = "engine-cuda"
+version = "4.0.2"
+dependencies = [
+ "cudarc",
+ "engine-cpu",
+ "log",
+ "metrics",
+ "pow-core",
+ "primitive-types 0.13.1",
+ "qp-poseidon-constants",
+]
+
 [[package]]
 name = "engine-gpu"
 version = "4.0.2"
@@ -597,6 +636,7 @@ dependencies = [
  "futures",
  "hex",
  "log",
+ "metrics",
  "pow-core",
  "primitive-types 0.13.1",
  "qp-plonky2",
@@ -636,7 +676,7 @@ source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "39cab71617ae0d63f51a36d69f866391735b51691dbda63cf6f96d042b63efeb"
 dependencies = [
  "libc",
- "windows-sys 0.52.0",
+ "windows-sys 0.61.2",
 ]
 
 [[package]]
@@ -1276,7 +1316,7 @@ checksum = "3640c1c38b8e4e43584d8df18be5fc6b0aa314ce6ebf51b53313d4306cca8e46"
 dependencies = [
  "hermit-abi",
  "libc",
- "windows-sys 0.52.0",
+ "windows-sys 0.61.2",
 ]
 
 [[package]]
@@ -1373,7 +1413,7 @@ source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "6aae1df220ece3c0ada96b8153459b67eebe9ae9212258bb0134ae60416fdf76"
 dependencies = [
  "libc",
- "libloading",
+ "libloading 0.8.9",
  "pkg-config",
 ]
 
@@ -1411,6 +1451,16 @@ dependencies = [
  "windows-link",
 ]
 
+[[package]]
+name = "libloading"
+version = "0.9.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "754ca22de805bb5744484a5b151a9e1a8e837d5dc232c2d7d8c2e3492edc8b60"
+dependencies = [
+ "cfg-if",
+ "windows-link",
+]
+
 [[package]]
 name = "libm"
 version = "0.2.16"
@@ -1527,6 +1577,7 @@ dependencies = [
  "anyhow",
  "crossbeam-channel",
  "engine-cpu",
+ "engine-cuda",
  "engine-gpu",
  "getrandom 0.2.17",
  "hex",
@@ -3702,7 +3753,7 @@ dependencies = [
  "js-sys",
  "khronos-egl",
  "libc",
- "libloading",
+ "libloading 0.8.9",
  "log",
  "metal",
  "naga",
@@ -3762,7 +3813,7 @@ version = "0.1.11"
 source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "c2a7b1c03c876122aa43f3020e6c3c3ee5c05081c9a00739faf7503aeba10d22"
 dependencies = [
- "windows-sys 0.48.0",
+ "windows-sys 0.61.2",
 ]
 
 [[package]]

Cargo.toml

@@ -1,6 +1,8 @@
 [workspace]
 members = [
+	"crates/bench-harness", # lair: quantus/miner#2
 	"crates/engine-cpu",
+	"crates/engine-cuda", # lair: quantus/miner#3
 	"crates/engine-gpu",
 	"crates/metrics",
 	"crates/miner-cli",

crates/bench-harness/Cargo.toml

@@ -0,0 +1,23 @@
+[package]
+name = "bench-harness"
+version.workspace = true
+edition.workspace = true
+publish = false
+description = "lair: reproducible GPU hashrate and parity measurement for the Quantus miner (quantus/miner#2)"
+
+[[bin]]
+name = "quantus-bench"
+path = "src/main.rs"
+
+[dependencies]
+engine-cpu = { path = "../engine-cpu" }
+engine-gpu = { path = "../engine-gpu" }
+engine-cuda = { path = "../engine-cuda" }
+pow-core = { path = "../pow-core" }
+primitive-types = { workspace = true }
+clap = { workspace = true, features = ["derive", "env"] }
+serde = { workspace = true }
+serde_json = { workspace = true, features = ["std"] }
+rand = { workspace = true }
+env_logger = { workspace = true }
+log = { workspace = true }

crates/bench-harness/bench-on-host.sh

@@ -0,0 +1,60 @@
+#!/usr/bin/env bash
+# lair: run quantus-bench on a mining host, as gitea_ci, with the miner paused.
+#
+# Executed over ssh by .gitea/workflows/bench.yaml. Everything that must be
+# serialised per host lives here under one flock: Gitea's workflow concurrency
+# group did NOT serialise two runs on the same host (observed: the second run's
+# scp hit ETXTBSY on the binary the first was executing), so the host is the
+# arbiter, not the forge.
+#
+# usage: bench-on-host.sh <binary> <duration_secs> <runs> <batch_size> <workers> <label> <out_json>
+set -euo pipefail
+
+BIN="$1"; DURATION="$2"; RUNS="$3"; BATCH="$4"; WORKERS="$5"; LABEL="$6"; OUT="$7"
+DIR=/var/lib/gitea_ci/bench
+LOCK="$DIR/.lock"
+
+exec 9>"$LOCK"
+if ! flock -w 1800 9; then
+  echo "another measurement has held $LOCK for 30 minutes; giving up" >&2
+  exit 1
+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
+
+echo "--- pause miner ---"
+was_active=0
+if systemctl is-active --quiet quantus-miner.service; then
+  was_active=1
+  sudo systemctl stop quantus-miner.service
+else
+  echo "quantus-miner.service was not active; nothing to pause"
+fi
+resume() {
+  if [ "$was_active" = 1 ]; then
+    echo "--- resume miner ---"
+    sudo systemctl start quantus-miner.service
+    systemctl is-active quantus-miner.service
+  fi
+}
+trap resume EXIT
+
+# Refuse to measure a card something else is using (on beast that would be
+# inference). A few seconds for the miner to release it.
+sleep 3
+util=$(nvidia-smi --query-gpu=utilization.gpu --format=csv,noheader,nounits | awk 'NR==1{printf "%d", $1}')
+echo "utilisation before measuring: ${util}%"
+if [ "$util" -gt 5 ]; then
+  echo "GPU is busy (${util}%) with the miner stopped; refusing to measure" >&2
+  exit 1
+fi
+
+echo "--- measure ---"
+cd "$DIR"
+RUST_LOG=info "$BIN" \
+  --duration-secs "$DURATION" --runs "$RUNS" --batch-size "$BATCH" --workers "$WORKERS" \
+  --expect-power-limit "$limit" --label "$LABEL" --json "$OUT"

crates/bench-harness/build.rs

@@ -0,0 +1,51 @@
+// lair: embed the git commit in the binary so `--version` and the build-info
+// metric identify a deployed build by commit, not by the workspace semver
+// (which does not change between commits on a branch that deploys on push).
+//
+// Resolution order:
+//   1. MINER_BUILD_SHA in the environment (CI sets it from the checked-out ref)
+//   2. `git rev-parse HEAD` of the workspace, with "-dirty" if the tree differs
+//   3. "unknown"
+use std::process::Command;
+
+fn git(args: &[&str]) -> Option<String> {
+    let out = Command::new("git").args(args).output().ok()?;
+    if !out.status.success() {
+        return None;
+    }
+    let s = String::from_utf8(out.stdout).ok()?;
+    let s = s.trim();
+    if s.is_empty() {
+        None
+    } else {
+        Some(s.to_string())
+    }
+}
+
+fn main() {
+    println!("cargo:rerun-if-env-changed=MINER_BUILD_SHA");
+
+    let sha = match std::env::var("MINER_BUILD_SHA") {
+        Ok(s) if !s.trim().is_empty() => s.trim().to_string(),
+        _ => match git(&["rev-parse", "--short=12", "HEAD"]) {
+            Some(head) => {
+                // Re-run when HEAD moves so a rebuild after a commit picks it up.
+                if let Some(dir) = git(&["rev-parse", "--git-dir"]) {
+                    println!("cargo:rerun-if-changed={dir}/HEAD");
+                    println!("cargo:rerun-if-changed={dir}/refs/heads");
+                }
+                let dirty = git(&["status", "--porcelain", "--untracked-files=no"])
+                    .map(|s| !s.is_empty())
+                    .unwrap_or(false);
+                if dirty {
+                    format!("{head}-dirty")
+                } else {
+                    head
+                }
+            }
+            None => "unknown".to_string(),
+        },
+    };
+
+    println!("cargo:rustc-env=MINER_BUILD_SHA={sha}");
+}

crates/bench-harness/src/main.rs

@@ -0,0 +1,506 @@
+//! lair: `quantus-bench`, the measurement gate for quantus/miner#2.
+//!
+//! Drives a GPU engine through the same `MinerEngine` trait the miner uses,
+//! records hashrate over fixed windows with the GPU's power and clock state
+//! captured before and after, verifies GPU/CPU parity on random jobs, and
+//! writes one JSON record so runs are comparable across commits.
+//!
+//! It deliberately measures one card per worker thread, with no node, no
+//! QUIC and no job churn: this is the kernel-plus-submission number. The
+//! production number, with all of that included, is quantus/miner#9.
+
+use clap::Parser;
+use engine_cpu::{AtomicBoolCancelCheck, EngineStatus, MinerEngine, Range};
+use engine_cuda::CudaEngine;
+use engine_gpu::GpuEngine;
+use primitive_types::U512;
+use rand::RngCore;
+use serde::Serialize;
+use std::process::Command;
+use std::sync::atomic::{AtomicBool, Ordering};
+use std::sync::Arc;
+use std::time::{Duration, Instant, SystemTime, UNIX_EPOCH};
+
+/// Reproducible GPU hashrate and parity measurement.
+#[derive(Parser, Debug)]
+#[command(version = VERSION, about)]
+struct Args {
+    /// Seconds per timed window.
+    #[arg(long, default_value_t = 30)]
+    duration_secs: u64,
+
+    /// Number of timed windows (median and spread are reported).
+    #[arg(long, default_value_t = 5)]
+    runs: usize,
+
+    /// Seconds of untimed warm-up before the first window.
+    #[arg(long, default_value_t = 5)]
+    warmup_secs: u64,
+
+    /// Nonces per GPU batch (the miner's --gpu-batch-size).
+    #[arg(long, default_value_t = 1_000_000)]
+    batch_size: u32,
+
+    /// GPU engine: auto (cuda if this binary carries a kernel and a driver is
+    /// present, else wgpu), cuda, or wgpu.
+    #[arg(long, default_value = "auto")]
+    engine: String,
+
+    /// Worker threads. The engine assigns threads to devices round-robin, so
+    /// on a multi-card host N threads measure N cards. More threads than cards
+    /// is allowed on purpose: two threads on one card overlap one thread's
+    /// readback with the other's dispatch, which is the submission bubble
+    /// quantus/miner#4 and #5 are about, measured without touching the miner.
+    #[arg(long, default_value_t = 1)]
+    workers: usize,
+
+    /// Random jobs for the GPU/CPU parity check (0 to skip).
+    #[arg(long, default_value_t = 25)]
+    parity_jobs: usize,
+
+    /// Refuse to run unless every GPU reports exactly this enforced power
+    /// limit in watts. Power limit is the largest confound on these cards.
+    #[arg(long)]
+    expect_power_limit: Option<u32>,
+
+    /// Commit the binary under test was built from.
+    #[arg(long, default_value = BUILD_SHA)]
+    commit: String,
+
+    /// Free-form label stored in the record (e.g. the PR number).
+    #[arg(long, default_value = "")]
+    label: String,
+
+    /// Write the JSON record here.
+    #[arg(long)]
+    json: Option<std::path::PathBuf>,
+}
+
+// Set by build.rs (same resolution as miner-cli's).
+const BUILD_SHA: &str = env!("MINER_BUILD_SHA");
+const VERSION: &str = concat!(
+    env!("CARGO_PKG_VERSION"),
+    " (",
+    env!("MINER_BUILD_SHA"),
+    ")"
+);
+
+#[derive(Serialize, Debug, Clone)]
+struct GpuState {
+    index: u32,
+    name: String,
+    driver: String,
+    power_limit_w: f64,
+    power_draw_w: f64,
+    sm_clock_mhz: f64,
+    /// The memory clock is a confound like the power limit: the miner never
+    /// touches VRAM, and a card with memory locked low gives the SMs more of
+    /// the same power budget (lair/quantus#10).
+    mem_clock_mhz: f64,
+    temperature_c: f64,
+}
+
+#[derive(Serialize, Debug)]
+struct WorkerResult {
+    worker: usize,
+    /// MH/s per timed window, in order.
+    windows_mhs: Vec<f64>,
+    median_mhs: f64,
+    /// (max - min) / median over the windows.
+    spread: f64,
+}
+
+#[derive(Serialize, Debug)]
+struct Parity {
+    jobs: usize,
+    found: usize,
+    ok: bool,
+}
+
+#[derive(Serialize, Debug)]
+struct Record {
+    schema: u32,
+    timestamp_unix: u64,
+    host: String,
+    commit: String,
+    label: String,
+    engine: String,
+    batch_size: u32,
+    duration_secs: u64,
+    runs: usize,
+    warmup_secs: u64,
+    workers: usize,
+    gpu_before: Vec<GpuState>,
+    gpu_after: Vec<GpuState>,
+    results: Vec<WorkerResult>,
+    /// Sum of worker medians: the host's number.
+    total_median_mhs: f64,
+    parity: Option<Parity>,
+}
+
+/// Either GPU engine behind the same trait; the harness drives them identically.
+enum Engine {
+    Wgpu(GpuEngine),
+    Cuda(CudaEngine),
+}
+
+impl Engine {
+    fn open(kind: &str, batch_size: u32) -> Engine {
+        match kind {
+            "cuda" => {
+                Engine::Cuda(CudaEngine::try_new(batch_size, 0).expect("CUDA engine init failed"))
+            }
+            "wgpu" => Engine::Wgpu(
+                GpuEngine::try_new(batch_size, 0, false).expect("GPU engine init failed"),
+            ),
+            "auto" => match CudaEngine::try_new(batch_size, 0) {
+                Ok(e) => Engine::Cuda(e),
+                Err(e) => {
+                    log::info!("CUDA engine unavailable ({e}); using wgpu");
+                    Engine::Wgpu(
+                        GpuEngine::try_new(batch_size, 0, false).expect("GPU engine init failed"),
+                    )
+                }
+            },
+            other => panic!("unknown --engine {other}; use auto, cuda or wgpu"),
+        }
+    }
+    fn device_count(&self) -> usize {
+        match self {
+            Engine::Wgpu(e) => e.device_count(),
+            Engine::Cuda(e) => e.device_count(),
+        }
+    }
+    fn as_dyn(&self) -> &dyn MinerEngine {
+        match self {
+            Engine::Wgpu(e) => e,
+            Engine::Cuda(e) => e,
+        }
+    }
+}
+
+fn nvidia_smi() -> Vec<GpuState> {
+    let out = Command::new("nvidia-smi")
+        .args([
+            "--query-gpu=index,name,driver_version,power.limit,power.draw,clocks.sm,clocks.mem,temperature.gpu",
+            "--format=csv,noheader,nounits",
+        ])
+        .output();
+    let Ok(out) = out else {
+        log::warn!("nvidia-smi not available; GPU state not recorded");
+        return Vec::new();
+    };
+    if !out.status.success() {
+        log::warn!("nvidia-smi failed; GPU state not recorded");
+        return Vec::new();
+    }
+    String::from_utf8_lossy(&out.stdout)
+        .lines()
+        .filter_map(|l| {
+            let f: Vec<&str> = l.split(',').map(str::trim).collect();
+            if f.len() != 8 {
+                return None;
+            }
+            let num = |s: &str| s.parse::<f64>().unwrap_or(f64::NAN);
+            Some(GpuState {
+                index: f[0].parse().ok()?,
+                name: f[1].to_string(),
+                driver: f[2].to_string(),
+                power_limit_w: num(f[3]),
+                power_draw_w: num(f[4]),
+                sm_clock_mhz: num(f[5]),
+                mem_clock_mhz: num(f[6]),
+                temperature_c: num(f[7]),
+            })
+        })
+        .collect()
+}
+
+fn hostname() -> String {
+    Command::new("hostname")
+        .arg("-f")
+        .output()
+        .ok()
+        .filter(|o| o.status.success())
+        .map(|o| String::from_utf8_lossy(&o.stdout).trim().to_string())
+        .unwrap_or_else(|| "unknown".into())
+}
+
+fn median(xs: &[f64]) -> f64 {
+    let mut v = xs.to_vec();
+    v.sort_by(|a, b| a.partial_cmp(b).unwrap());
+    let n = v.len();
+    if n == 0 {
+        return f64::NAN;
+    }
+    if n % 2 == 1 {
+        v[n / 2]
+    } else {
+        (v[n / 2 - 1] + v[n / 2]) / 2.0
+    }
+}
+
+/// One search window: search from a random start until `cancel` fires,
+/// returning hashes and elapsed seconds. The difficulty is high enough that
+/// a solution is effectively impossible, matching production where the second
+/// squeeze is almost never taken.
+fn window(engine: &dyn MinerEngine, secs: u64, rng: &mut impl RngCore) -> (u64, f64) {
+    let mut header = [0u8; 32];
+    rng.fill_bytes(&mut header);
+    let ctx = engine.prepare_context(header, U512::from(1u64) << 200);
+
+    let mut start_bytes = [0u8; 64];
+    rng.fill_bytes(&mut start_bytes);
+    start_bytes[0] = 0; // keep clear of the top so the range cannot wrap
+    let start = U512::from_big_endian(&start_bytes);
+    let range = Range {
+        start,
+        end: start + (U512::from(1u64) << 200),
+    };
+
+    let flag = Arc::new(AtomicBool::new(false));
+    let timer_flag = flag.clone();
+    let timer = std::thread::spawn(move || {
+        std::thread::sleep(Duration::from_secs(secs));
+        timer_flag.store(true, Ordering::Relaxed);
+    });
+
+    let t0 = Instant::now();
+    let status = engine.search_range(&ctx, range, &AtomicBoolCancelCheck(&flag));
+    let elapsed = t0.elapsed().as_secs_f64();
+    timer.join().expect("timer thread");
+
+    let hashes = match status {
+        EngineStatus::Cancelled { hash_count } | EngineStatus::Exhausted { hash_count } => {
+            hash_count
+        }
+        EngineStatus::Found { hash_count, .. } => {
+            log::warn!("window found a solution at 2^-200 odds; counting hashes anyway");
+            hash_count
+        }
+        other => panic!("unexpected engine status: {other:?}"),
+    };
+    (hashes, elapsed)
+}
+
+fn parity(engine: &dyn MinerEngine, jobs: usize) -> Parity {
+    // Same shape as engine-gpu's gpu_cpu_parity example, kept in lockstep with it.
+    let mut rng = rand::rng();
+    let cancel_flag = AtomicBool::new(false);
+    let cancel = AtomicBoolCancelCheck(&cancel_flag);
+    let mut found = 0usize;
+    let mut ok = true;
+    for job in 0..jobs {
+        let mut header = [0u8; 32];
+        rng.fill_bytes(&mut header);
+        let ctx = engine.prepare_context(header, U512::from(100_000u64));
+        let start = if job == 0 {
+            // Cross a 2^256 boundary to exercise the midstate batch clamp.
+            (U512::from(3u64) << 256) - U512::from(1_000u64)
+        } else {
+            let mut b = [0u8; 64];
+            rng.fill_bytes(&mut b);
+            b[0] = 0;
+            U512::from_big_endian(&b)
+        };
+        let range = Range {
+            start,
+            end: start + U512::from(10_000_000u64),
+        };
+        match engine.search_range(&ctx, range.clone(), &cancel) {
+            EngineStatus::Found { candidate, .. } => {
+                let cpu = pow_core::hash_from_nonce(&ctx, candidate.nonce);
+                let in_range = candidate.nonce >= range.start && candidate.nonce <= range.end;
+                if cpu != candidate.hash || cpu >= ctx.target || !in_range {
+                    log::error!(
+                        "parity job {job}: nonce {} gpu_hash {} cpu_hash {} in_range {in_range}",
+                        candidate.nonce,
+                        candidate.hash,
+                        cpu
+                    );
+                    ok = false;
+                }
+                found += 1;
+            }
+            EngineStatus::Exhausted { .. } => {}
+            other => {
+                log::error!("parity job {job}: unexpected status {other:?}");
+                ok = false;
+            }
+        }
+    }
+    if found == 0 {
+        log::error!("parity: no solutions found across {jobs} jobs");
+        ok = false;
+    }
+    Parity { jobs, found, ok }
+}
+
+fn main() {
+    env_logger::Builder::from_env(env_logger::Env::default().default_filter_or("info")).init();
+    let args = Args::parse();
+    if args.runs == 0 || args.workers == 0 {
+        eprintln!("--runs and --workers must be at least 1");
+        std::process::exit(2);
+    }
+
+    let gpu_before = nvidia_smi();
+    if let Some(want) = args.expect_power_limit {
+        for g in &gpu_before {
+            if (g.power_limit_w - want as f64).abs() > 0.5 {
+                eprintln!(
+                    "gpu {} ({}) power limit is {:.0} W, expected {want} W; refusing to measure",
+                    g.index, g.name, g.power_limit_w
+                );
+                std::process::exit(3);
+            }
+        }
+    }
+
+    let engine = Arc::new(Engine::open(&args.engine, args.batch_size));
+    let devices = engine.device_count();
+    if args.workers > devices {
+        log::warn!(
+            "{} workers on {devices} device(s): threads share cards, so per-worker numbers \
+             are shares and total_median_mhs is the number to compare",
+            args.workers
+        );
+    }
+    log::info!(
+        "engine {} with {devices} device(s); measuring {} worker(s), {} x {}s windows after {}s warm-up, batch {}",
+        engine.as_dyn().name(),
+        args.workers,
+        args.runs,
+        args.duration_secs,
+        args.warmup_secs,
+        args.batch_size
+    );
+
+    let mut handles = Vec::new();
+    for w in 0..args.workers {
+        let engine = engine.clone();
+        let (runs, dur, warm) = (args.runs, args.duration_secs, args.warmup_secs);
+        handles.push(std::thread::spawn(move || {
+            let mut rng = rand::rng();
+            if warm > 0 {
+                let _ = window(engine.as_dyn(), warm, &mut rng);
+            }
+            let mut windows = Vec::with_capacity(runs);
+            for i in 0..runs {
+                let (h, s) = window(engine.as_dyn(), dur, &mut rng);
+                let mhs = h as f64 / s / 1e6;
+                log::info!("worker {w} window {i}: {h} hashes in {s:.2}s = {mhs:.2} MH/s");
+                windows.push(mhs);
+            }
+            let med = median(&windows);
+            let (mn, mx) = windows
+                .iter()
+                .fold((f64::INFINITY, f64::NEG_INFINITY), |(a, b), &x| {
+                    (a.min(x), b.max(x))
+                });
+            WorkerResult {
+                worker: w,
+                windows_mhs: windows,
+                median_mhs: med,
+                spread: if med > 0.0 { (mx - mn) / med } else { f64::NAN },
+            }
+        }));
+    }
+    let results: Vec<WorkerResult> = handles
+        .into_iter()
+        .map(|h| h.join().expect("worker thread"))
+        .collect();
+    let gpu_after = nvidia_smi();
+
+    let parity_result = if args.parity_jobs > 0 {
+        Some(parity(engine.as_dyn(), args.parity_jobs))
+    } else {
+        None
+    };
+
+    let record = Record {
+        schema: 2,
+        timestamp_unix: SystemTime::now()
+            .duration_since(UNIX_EPOCH)
+            .map(|d| d.as_secs())
+            .unwrap_or(0),
+        host: hostname(),
+        commit: args.commit.clone(),
+        label: args.label.clone(),
+        engine: engine.as_dyn().name().to_string(),
+        batch_size: args.batch_size,
+        duration_secs: args.duration_secs,
+        runs: args.runs,
+        warmup_secs: args.warmup_secs,
+        workers: args.workers,
+        gpu_before,
+        gpu_after,
+        total_median_mhs: results.iter().map(|r| r.median_mhs).sum(),
+        results,
+        parity: parity_result,
+    };
+
+    // Human summary on stdout (the workflow appends it to the job summary).
+    println!("## quantus-bench {} on {}", record.commit, record.host);
+    println!();
+    println!("| worker | median MH/s | spread | windows |");
+    println!("| --- | --- | --- | --- |");
+    for r in &record.results {
+        let w: Vec<String> = r.windows_mhs.iter().map(|x| format!("{x:.1}")).collect();
+        println!(
+            "| {} | {:.2} | {:.1}% | {} |",
+            r.worker,
+            r.median_mhs,
+            r.spread * 100.0,
+            w.join(", ")
+        );
+    }
+    println!();
+    println!(
+        "total median: **{:.2} MH/s** (batch {}, {} x {}s, engine {})",
+        record.total_median_mhs,
+        record.batch_size,
+        record.runs,
+        record.duration_secs,
+        record.engine
+    );
+    for g in &record.gpu_after {
+        println!(
+            "- gpu {} {}: driver {}, limit {:.0} W, draw {:.0} W, sm {:.0} MHz, mem {:.0} MHz, {:.0} C",
+            g.index,
+            g.name,
+            g.driver,
+            g.power_limit_w,
+            g.power_draw_w,
+            g.sm_clock_mhz,
+            g.mem_clock_mhz,
+            g.temperature_c
+        );
+    }
+    if let Some(p) = &record.parity {
+        println!(
+            "- parity: {} ({}/{} jobs found solutions, all verified against CPU)",
+            if p.ok { "OK" } else { "FAILED" },
+            p.found,
+            p.jobs
+        );
+    }
+
+    if let Some(path) = &args.json {
+        let s = serde_json::to_string_pretty(&record).expect("serialise record");
+        std::fs::write(path, s).expect("write json record");
+        log::info!("record written to {}", path.display());
+    }
+
+    let spread_bad = record.results.iter().any(|r| r.spread > 0.02);
+    if spread_bad {
+        eprintln!("spread above 2% on at least one worker; run is not stable enough to compare");
+    }
+    if record.parity.as_ref().is_some_and(|p| !p.ok) {
+        eprintln!("PARITY FAILED");
+        std::process::exit(4);
+    }
+    if spread_bad {
+        std::process::exit(5);
+    }
+}

crates/engine-cuda/Cargo.toml

@@ -0,0 +1,17 @@
+[package]
+name = "engine-cuda"
+version.workspace = true
+edition.workspace = true
+publish = false
+description = "lair: native CUDA mining engine behind MinerEngine (quantus/miner#3)"
+
+[dependencies]
+engine-cpu = { path = "../engine-cpu" }
+pow-core = { path = "../pow-core" }
+metrics = { path = "../metrics" }
+primitive-types = { workspace = true }
+log = { workspace = true }
+cudarc = { version = "0.19", default-features = false, features = ["driver", "nvrtc", "dynamic-loading", "cuda-13000"] }
+
+[build-dependencies]
+qp-poseidon-constants = "1.1"

crates/engine-cuda/build.rs

@@ -0,0 +1,196 @@
+//! lair: build the CUDA mining kernel into a fat binary (quantus/miner#3, #8).
+//!
+//! 1. Generate `poseidon2_constants.cuh` from `qp-poseidon-constants`, so a
+//!    change to the hash at origin reaches the kernel as a dependency bump and
+//!    never as a hand-copied table (rule from quantus/miner#1).
+//! 2. If `nvcc` is available, compile `src/kernels/mining.cu` once with one
+//!    cubin per fleet architecture plus PTX for the newest, into
+//!    `$OUT_DIR/mining.fatbin`. The driver picks the matching cubin at load;
+//!    an unknown future card JIT-compiles the PTX.
+//! 3. If `nvcc` is not available (the `rust` lint runner, a workstation without
+//!    CUDA), write an empty fat binary and warn. The crate still compiles;
+//!    `CudaEngine::try_new` then fails with a clear message and the miner
+//!    falls back to the wgpu engine.
+//!
+//! Knobs (environment):
+//!   NVCC / CUDA_HOME                     where nvcc is
+//!   MINER_CUDA_ARCHS                     comma-separated SM list, default 86,89,120
+//!   MINER_NVCC_CCBIN                     host compiler for nvcc (-ccbin)
+//!   MINER_CUDA_ALLOW_UNSUPPORTED_COMPILER=1  pass -allow-unsupported-compiler
+//!   MINER_CUDA_REQUIRE=1                 fail the build instead of warning when nvcc is missing
+//!   MINER_NVCC_FLAGS                     extra nvcc arguments (whitespace-separated), e.g. -DLAIR_LAZY_ADD=0
+
+use std::env;
+use std::fs;
+use std::path::{Path, PathBuf};
+use std::process::Command;
+
+fn write_constants(out: &Path) {
+    use qp_poseidon_constants as c;
+    let mut s = String::new();
+    s.push_str("// Generated by engine-cuda/build.rs from qp-poseidon-constants. Do not edit.\n");
+    s.push_str("#pragma once\n");
+    let arr = |name: &str, v: &[u64]| {
+        let body: Vec<String> = v.iter().map(|x| format!("0x{x:016x}ull")).collect();
+        format!(
+            "__device__ __constant__ unsigned long long {}[{}] = {{\n    {}\n}};\n",
+            name,
+            v.len(),
+            body.join(",\n    ")
+        )
+    };
+    let arr2 = |name: &str, v: &[[u64; 12]]| {
+        let rows: Vec<String> = v
+            .iter()
+            .map(|r| {
+                let body: Vec<String> = r.iter().map(|x| format!("0x{x:016x}ull")).collect();
+                format!("    {{{}}}", body.join(", "))
+            })
+            .collect();
+        format!(
+            "__device__ __constant__ unsigned long long {}[{}][12] = {{\n{}\n}};\n",
+            name,
+            v.len(),
+            rows.join(",\n")
+        )
+    };
+    s.push_str(&arr("RC_INTERNAL", &c::POSEIDON2_INTERNAL_CONSTANTS_RAW));
+    s.push_str(&arr2(
+        "RC_INITIAL",
+        &c::POSEIDON2_INITIAL_EXTERNAL_CONSTANTS_RAW,
+    ));
+    s.push_str(&arr2(
+        "RC_TERMINAL",
+        &c::POSEIDON2_TERMINAL_EXTERNAL_CONSTANTS_RAW,
+    ));
+    s.push_str(&arr("MDS_DIAG", &c::POSEIDON2_MATRIX_DIAG_12_RAW));
+    s.push_str(&format!(
+        "#define SPONGE_WIDTH {}\n#define N_EXTERNAL_HALF {}\n#define N_INTERNAL {}\n",
+        c::SPONGE_WIDTH,
+        c::POSEIDON2_EXTERNAL_ROUNDS / 2,
+        c::POSEIDON2_INTERNAL_ROUNDS
+    ));
+    fs::write(out.join("poseidon2_constants.cuh"), s).expect("write constants header");
+}
+
+fn find_nvcc() -> Option<PathBuf> {
+    if let Ok(p) = env::var("NVCC") {
+        let p = PathBuf::from(p);
+        if p.is_file() {
+            return Some(p);
+        }
+    }
+    if let Ok(home) = env::var("CUDA_HOME").or_else(|_| env::var("CUDA_PATH")) {
+        let p = PathBuf::from(home).join("bin").join("nvcc");
+        if p.is_file() {
+            return Some(p);
+        }
+    }
+    if let Ok(path) = env::var("PATH") {
+        for dir in env::split_paths(&path) {
+            let p = dir.join("nvcc");
+            if p.is_file() {
+                return Some(p);
+            }
+        }
+    }
+    // The toolkit the fleet drivers support; prefer an exact 13.0 over "latest".
+    for cand in [
+        "/usr/local/cuda-13.0",
+        "/usr/local/cuda-13",
+        "/usr/local/cuda",
+    ] {
+        let p = PathBuf::from(cand).join("bin").join("nvcc");
+        if p.is_file() {
+            return Some(p);
+        }
+    }
+    None
+}
+
+fn main() {
+    println!("cargo:rerun-if-changed=build.rs");
+    println!("cargo:rerun-if-changed=src/kernels/mining.cu");
+    for v in [
+        "NVCC",
+        "CUDA_HOME",
+        "CUDA_PATH",
+        "MINER_CUDA_ARCHS",
+        "MINER_NVCC_CCBIN",
+        "MINER_CUDA_ALLOW_UNSUPPORTED_COMPILER",
+        "MINER_CUDA_REQUIRE",
+        "MINER_NVCC_FLAGS",
+    ] {
+        println!("cargo:rerun-if-env-changed={v}");
+    }
+
+    let out = PathBuf::from(env::var("OUT_DIR").expect("OUT_DIR"));
+    write_constants(&out);
+
+    let archs = env::var("MINER_CUDA_ARCHS").unwrap_or_else(|_| "86,89,120".to_string());
+    let sms: Vec<String> = archs
+        .split(',')
+        .map(|s| s.trim().trim_start_matches("sm_").to_string())
+        .filter(|s| !s.is_empty())
+        .collect();
+    println!("cargo:rustc-env=MINER_CUDA_ARCHS={}", sms.join(","));
+
+    let fatbin = out.join("mining.fatbin");
+    let Some(nvcc) = find_nvcc() else {
+        if env::var("MINER_CUDA_REQUIRE")
+            .map(|v| v == "1")
+            .unwrap_or(false)
+        {
+            panic!("nvcc not found and MINER_CUDA_REQUIRE=1; set NVCC or CUDA_HOME");
+        }
+        println!("cargo:warning=engine-cuda: nvcc not found; CUDA kernel not compiled into this binary (wgpu engine will be used)");
+        fs::write(&fatbin, []).expect("write empty fatbin");
+        return;
+    };
+
+    let mut cmd = Command::new(&nvcc);
+    cmd.arg("-fatbin")
+        .arg("-O3")
+        .arg("-std=c++17")
+        .arg("-I")
+        .arg(&out)
+        .arg("-o")
+        .arg(&fatbin);
+    if let Ok(cc) = env::var("MINER_NVCC_CCBIN") {
+        cmd.arg("-ccbin").arg(cc);
+    }
+    if env::var("MINER_CUDA_ALLOW_UNSUPPORTED_COMPILER")
+        .map(|v| v == "1")
+        .unwrap_or(false)
+    {
+        cmd.arg("-allow-unsupported-compiler");
+    }
+    if let Ok(extra) = env::var("MINER_NVCC_FLAGS") {
+        for a in extra.split_whitespace() {
+            cmd.arg(a);
+        }
+    }
+    for sm in &sms {
+        cmd.arg("--generate-code")
+            .arg(format!("arch=compute_{sm},code=sm_{sm}"));
+    }
+    if let Some(newest) = sms.iter().max_by_key(|s| s.parse::<u32>().unwrap_or(0)) {
+        cmd.arg("--generate-code")
+            .arg(format!("arch=compute_{newest},code=compute_{newest}"));
+    }
+    cmd.arg("src/kernels/mining.cu");
+
+    println!(
+        "cargo:warning=engine-cuda: {} for sm_{{{}}}",
+        nvcc.display(),
+        sms.join(",")
+    );
+    let status = cmd.status().expect("run nvcc");
+    if !status.success() {
+        panic!("nvcc failed ({status}); see output above");
+    }
+    let size = fs::metadata(&fatbin).map(|m| m.len()).unwrap_or(0);
+    if size == 0 {
+        panic!("nvcc produced an empty fat binary");
+    }
+}

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

@@ -0,0 +1,539 @@
+// lair: Poseidon2-over-Goldilocks mining kernel, native CUDA (quantus/miner#3).
+//
+// Same host contract as engine-gpu's mining_u64.wgsl, and bit-exact with it and
+// with pow_core:
+//   - the host precomputes the sponge midstate after absorbing the 32-byte
+//     header and the high 32 bytes of the big-endian nonce (pow_core::
+//     mining_midstate), so each nonce costs 2 permutations instead of 5;
+//   - a batch never carries into the high 256 bits of the nonce, so only the
+//     low 8 u32 limbs are incremented here;
+//   - the first squeeze yields the most significant 256 bits of the hash, which
+//     decide hash-vs-target on their own unless they exactly equal the target's
+//     high half; only such candidates pay for the second squeeze.
+//
+// What CUDA buys over WGSL: a 64x64 -> 128-bit multiply is `__umul64hi` plus a
+// plain multiply, instead of four 32-bit partial products with carry
+// reconstruction. The field multiply is the whole cost of this hash.
+//
+// Field elements are kept in lazy (non-canonical) form below 2^64, exactly as
+// the WGSL kernel does, and canonicalised only when bytes are produced.
+
+#include <stdint.h>
+#include "poseidon2_constants.cuh"
+
+typedef unsigned long long u64;
+typedef unsigned int u32;
+
+#define P64 0xFFFFFFFF00000001ull
+// 2^64 mod P = 2^32 - 1
+#define EPS 0xFFFFFFFFull
+
+// lair: whether gf_add and gf_reduce should use inline-PTX carry arithmetic or
+// plain C is architecture-dependent, and the difference is large in both
+// directions. An `asm` block is opaque to nvcc's optimiser, so carry arithmetic
+// written as PTX blocks common-subexpression elimination and strength reduction
+// across neighbouring field operations. On sm_120 that costs a quarter of the
+// kernel; on sm_86/89 the hand-written sequence still wins.
+//
+// Whole-kernel instructions / ALU ops (ptxas + cuobjdump, CUDA 13.0) and
+// measured hashrate, three interleaved rounds each, parity 40/40 vs CPU:
+//
+//              sm_120                         sm_86/89
+//   CARRY=1    37,432 / 26,688   1126.6      37,027   benjy 461.6  quadbrat 85.5
+//   CARRY=0    27,837 / 19,519   1333.3      41,196   benjy 401.5  quadbrat 76.3
+//
+// So +18.3% on the 5090s and -13.0% / -10.8% on the 4090 and 3060. Hence the
+// per-arch default. Unlike LAIR_INT_UNROLL -- where the static measures pointed
+// one way and the cards' power envelope the other -- this one was measured on
+// hardware on all three architectures, and the static count predicted the sign
+// correctly on each.
+//
+// LAIR_PTX_ACC stays 1 everywhere: in the Acc accumulators the add-with-carry
+// pair is the whole operation, so there is nothing around it to optimise. On
+// sm_120, carry off with Acc off measured 1285.8 against 1333.3 with Acc on.
+//
+// PR #17 measured the PTX carry path as a win on every card, and that was true
+// of the kernel it was written for -- which spilled 104 bytes per thread. Once
+// the spills went (see MiningUniforms) the trade reversed on Blackwell only.
+// Re-measure per architecture before touching this.
+#ifndef LAIR_PTX_CARRY
+#  if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 1200
+#    define LAIR_PTX_CARRY 0
+#  else
+#    define LAIR_PTX_CARRY 1
+#  endif
+#endif
+#ifndef LAIR_PTX_ACC
+#define LAIR_PTX_ACC 1
+#endif
+// Fold the next round's constant into the linear layer's final reduction
+// instead of a separate gf_add per element before the S-box.
+#ifndef LAIR_FOLD_RC
+#define LAIR_FOLD_RC 1
+#endif
+
+#if LAIR_PTX_CARRY
+// Carry-flag arithmetic: the wrap is read from the condition code instead of
+// a compare, and folded with a multiply-add (c * EPS is exact for c in {0,1}).
+__device__ __forceinline__ u64 gf_add(u64 a, u64 b) {
+    u64 s0, c1, s1, c2;
+    asm("add.cc.u64 %0, %2, %3;\n\taddc.u64 %1, 0, 0;" : "=l"(s0), "=l"(c1) : "l"(a), "l"(b));
+    asm("add.cc.u64 %0, %2, %3;\n\taddc.u64 %1, 0, 0;" : "=l"(s1), "=l"(c2) : "l"(s0), "l"(c1 * EPS));
+    return s1 + c2 * EPS;
+}
+
+__device__ __forceinline__ u64 gf_reduce(u64 lo, u64 hi) {
+    u64 hi_hi = hi >> 32;
+    u64 hi_lo = hi & EPS;
+    u64 t0, bm;
+    // bm = -borrow: all ones when lo < hi_hi.
+    asm("sub.cc.u64 %0, %2, %3;\n\tsubc.u64 %1, 0, 0;" : "=l"(t0), "=l"(bm) : "l"(lo), "l"(hi_hi));
+    t0 -= (bm & EPS);
+    u64 t1 = hi_lo * EPS;
+    u64 t2, c;
+    asm("add.cc.u64 %0, %2, %3;\n\taddc.u64 %1, 0, 0;" : "=l"(t2), "=l"(c) : "l"(t0), "l"(t1));
+    return t2 + c * EPS;
+}
+#else
+// a + b mod P, lazy. A wrapped carry folds back as 2^64 = EPS (mod P); the
+// second fold can only be needed when the first one wrapped.
+__device__ __forceinline__ u64 gf_add(u64 a, u64 b) {
+    u64 s0 = a + b;
+    u64 c1 = s0 < a;
+    u64 s1 = s0 + (c1 ? EPS : 0ull);
+    u64 c2 = c1 & (s1 < s0);
+    return s1 + (c2 ? EPS : 0ull);
+}
+
+// Reduce lo + hi * 2^64 mod P using 2^64 = EPS and 2^96 = -1 (mod P).
+__device__ __forceinline__ u64 gf_reduce(u64 lo, u64 hi) {
+    u64 hi_hi = hi >> 32;
+    u64 hi_lo = hi & EPS;
+    u64 t0 = lo - hi_hi;
+    if (lo < hi_hi) t0 -= EPS;
+    u64 t1 = hi_lo * EPS;
+    u64 t2 = t0 + t1;
+    if (t2 < t0) t2 += EPS;
+    return t2;
+}
+#endif
+
+__device__ __forceinline__ u64 gf_mul(u64 a, u64 b) {
+    return gf_reduce(a * b, __umul64hi(a, b));
+}
+
+// Squaring deliberately uses the general product. nvcc does not specialise
+// `a * a`: it emits the same six IMAD.WIDE in 48 instructions as gf_mul
+// (measured, sm_120). The textbook saving is real but does not pay here --
+// a = a1*2^32 + a0 gives a^2 = a0^2 + a0*a1*2^33 + a1^2*2^64, three 32x32
+// multiplies instead of four, and it was implemented, verified bit-exact
+// against __int128 over 40M values, and measured 5.7% SLOWER on beast's 5090s
+// (1072.9 vs 1137.7 MH/s, three interleaved rounds). It saves 6.7% of the
+// kernel's widening multiplies and costs 11% more instructions: 5 multiplies
+// in 56 instructions against 6 in 48. Four formulations (condition-code
+// carry, compare carry, and two mad.wide.u32 variants) all compiled to the
+// same 5/56, so 56 is the floor. The scarce resource in this kernel is
+// instruction issue, not multiply throughput -- which is also why raising
+// occupancy lost (see LAIR_INT_UNROLL). Anything that wins here has to remove
+// work, not re-express it.
+__device__ __forceinline__ u64 gf_sqr(u64 a) {
+    return gf_reduce(a * a, __umul64hi(a, a));
+}
+
+// x^7
+__device__ __forceinline__ u64 gf_sbox(u64 x) {
+    u64 x2 = gf_sqr(x);
+    u64 x4 = gf_sqr(x2);
+    u64 x6 = gf_mul(x4, x2);
+    return gf_mul(x6, x);
+}
+
+__device__ __forceinline__ u64 gf_canon(u64 a) {
+    return a - (a >= P64 ? P64 : 0ull);
+}
+
+#ifndef LAIR_LAZY_ADD
+#define LAIR_LAZY_ADD 1
+#endif
+// lair: 2 stays the default on every architecture. Unroll 1 looks better on
+// sm_120 by every static measure -- 78 registers instead of 106, no spill
+// either way, a third resident block per SM instead of two -- and measured
+// 0.92% SLOWER on beast's 5090s (1128.4 vs 1138.8 MH/s, three interleaved
+// rounds, spread <=0.1%). These cards are power-bound, not latency-bound: at
+// the 400 W cap the extra resident warps buy power draw, and the card clocks
+// down to pay for it (2325/2385 MHz against 2355/2415). Occupancy is not the
+// lever here; work per hash is. Do not "fix" this without a measurement.
+#ifndef LAIR_INT_UNROLL
+#define LAIR_INT_UNROLL 2
+#endif
+// nvcc does not macro-expand `#pragma unroll N`; stringise through _Pragma.
+#define LAIR_PRAGMA(x) _Pragma(#x)
+#define LAIR_UNROLL(n) LAIR_PRAGMA(unroll n)
+#ifndef LAIR_PTX_CARRY
+#  if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 1200
+#    define LAIR_PTX_CARRY 0
+#  else
+#    define LAIR_PTX_CARRY 1
+#  endif
+#endif
+#ifndef LAIR_PTX_ACC
+#define LAIR_PTX_ACC 1
+#endif
+// Fold the next round's constant into the linear layer's final reduction
+// instead of a separate gf_add per element before the S-box.
+#ifndef LAIR_FOLD_RC
+#define LAIR_FOLD_RC 1
+#endif
+#ifndef LAIR_NOINLINE_PERMUTE
+#define LAIR_NOINLINE_PERMUTE 0
+#endif
+
+// Deferred-carry accumulator: the value is lo + hi * 2^64 with hi small. Adds
+// cost an add-with-carry instead of two compare-and-fold steps; the folds are
+// paid once per output when the accumulator is reduced. Sound as long as
+// hi < 2^32, which a linear layer's sums (at most a few dozen terms) satisfy.
+struct Acc {
+    u64 lo;
+    u32 hi;
+};
+
+__device__ __forceinline__ Acc acc_of(u64 a) {
+    Acc r; r.lo = a; r.hi = 0u; return r;
+}
+
+#if LAIR_PTX_ACC
+__device__ __forceinline__ Acc acc_add(Acc a, u64 b) {
+    u64 s; u32 c;
+    asm("add.cc.u64 %0, %2, %3;\n\taddc.u32 %1, 0, 0;" : "=l"(s), "=r"(c) : "l"(a.lo), "l"(b));
+    a.lo = s;
+    a.hi += c;
+    return a;
+}
+
+__device__ __forceinline__ Acc acc_add2(Acc a, Acc b) {
+    u64 s; u32 c;
+    asm("add.cc.u64 %0, %2, %3;\n\taddc.u32 %1, %4, %5;" : "=l"(s), "=r"(c) : "l"(a.lo), "l"(b.lo), "r"(a.hi), "r"(b.hi));
+    a.lo = s;
+    a.hi = c;
+    return a;
+}
+#else
+__device__ __forceinline__ Acc acc_add(Acc a, u64 b) {
+    u64 s = a.lo + b;
+    a.hi += (s < a.lo) ? 1u : 0u;
+    a.lo = s;
+    return a;
+}
+
+__device__ __forceinline__ Acc acc_add2(Acc a, Acc b) {
+    u64 s = a.lo + b.lo;
+    a.hi += b.hi + ((s < a.lo) ? 1u : 0u);
+    a.lo = s;
+    return a;
+}
+#endif
+
+// lo + hi * 2^64 = lo + hi * EPS (mod P); hi * EPS < 2^64, so this is one
+// lazy gf_add.
+__device__ __forceinline__ u64 acc_reduce(Acc a) {
+    return gf_add(a.lo, (u64)a.hi * EPS);
+}
+
+// a * b + c mod P: the addend is folded into the 128-bit product before the
+// single reduction, saving a gf_add per term.
+__device__ __forceinline__ u64 gf_mul_add(u64 a, u64 b, u64 c) {
+    u64 lo = a * b;
+    u64 hi = __umul64hi(a, b);
+    u64 s = lo + c;
+    hi += (s < lo) ? 1ull : 0ull;
+    return gf_reduce(s, hi);
+}
+
+// a * b + (c.lo + c.hi * 2^64) mod P, for a deferred-carry addend.
+__device__ __forceinline__ u64 gf_mul_add_acc(u64 a, u64 b, Acc c) {
+    u64 lo = a * b;
+    u64 hi = __umul64hi(a, b);
+    u64 s = lo + c.lo;
+    hi += ((s < lo) ? 1ull : 0ull) + (u64)c.hi;
+    return gf_reduce(s, hi);
+}
+
+#if LAIR_LAZY_ADD
+// External linear layer with deferred carries: every output is a sum of a
+// handful of inputs, reduced once.
+// `rc` is the next round's constant row to fold into the outputs, or nullptr.
+__device__ __forceinline__ void ext_layer_rc(u64 st[12], const u64* rc) {
+    Acc acc[12];
+#pragma unroll
+    for (int chunk = 0; chunk < 3; chunk++) {
+        int o = chunk * 4;
+        u64 x0 = st[o], x1 = st[o + 1], x2 = st[o + 2], x3 = st[o + 3];
+        Acc t01 = acc_add(acc_of(x0), x1);
+        Acc t23 = acc_add(acc_of(x2), x3);
+        Acc t0123 = acc_add2(t01, t23);
+        Acc t01123 = acc_add(t0123, x1);
+        Acc t01233 = acc_add(t0123, x3);
+        acc[o + 3] = acc_add(acc_add(t01233, x0), x0);
+        acc[o + 1] = acc_add(acc_add(t01123, x2), x2);
+        acc[o] = acc_add2(t01123, t01);
+        acc[o + 2] = acc_add2(t01233, t23);
+    }
+    Acc sums[4];
+#pragma unroll
+    for (int k = 0; k < 4; k++) {
+        sums[k] = acc_add2(acc_add2(acc[k], acc[k + 4]), acc[k + 8]);
+    }
+#pragma unroll
+    for (int i = 0; i < 12; i++) {
+        Acc o = acc_add2(acc[i], sums[i & 3]);
+        if (rc != nullptr) o = acc_add(o, rc[i]);
+        st[i] = acc_reduce(o);
+    }
+}
+
+__device__ __forceinline__ void ext_layer(u64 st[12]) { ext_layer_rc(st, nullptr); }
+
+// 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(u64 st[12], u64 rc0, bool has_rc0, const u64* rc_row) {
+    Acc s = acc_of(st[0]);
+#pragma unroll
+    for (int i = 1; i < 12; i++) s = acc_add(s, st[i]);
+    u64 sum = acc_reduce(s);
+    if (rc_row != nullptr) {
+#pragma unroll
+        for (int i = 0; i < 12; i++) st[i] = gf_mul_add_acc(st[i], MDS_DIAG[i], acc_add(acc_of(sum), rc_row[i]));
+    } else {
+        if (has_rc0) {
+            st[0] = gf_mul_add_acc(st[0], MDS_DIAG[0], acc_add(acc_of(sum), rc0));
+        } else {
+            st[0] = gf_mul_add(st[0], MDS_DIAG[0], sum);
+        }
+#pragma unroll
+        for (int i = 1; i < 12; i++) st[i] = gf_mul_add(st[i], MDS_DIAG[i], sum);
+    }
+}
+
+__device__ __forceinline__ void int_layer(u64 st[12]) { int_layer_rc(st, 0ull, false, nullptr); }
+#else
+// External linear layer: 4x4 MDS on each chunk, then circulant sums.
+__device__ __forceinline__ void ext_layer(u64 st[12]) {
+#pragma unroll
+    for (int chunk = 0; chunk < 3; chunk++) {
+        int o = chunk * 4;
+        u64 x0 = st[o], x1 = st[o + 1], x2 = st[o + 2], x3 = st[o + 3];
+        u64 t01 = gf_add(x0, x1);
+        u64 t23 = gf_add(x2, x3);
+        u64 t0123 = gf_add(t01, t23);
+        u64 t01123 = gf_add(t0123, x1);
+        u64 t01233 = gf_add(t0123, x3);
+        st[o + 3] = gf_add(t01233, gf_add(x0, x0));
+        st[o + 1] = gf_add(t01123, gf_add(x2, x2));
+        st[o] = gf_add(t01123, t01);
+        st[o + 2] = gf_add(t01233, t23);
+    }
+    u64 sums[4];
+#pragma unroll
+    for (int k = 0; k < 4; k++) {
+        sums[k] = gf_add(gf_add(st[k], st[k + 4]), st[k + 8]);
+    }
+#pragma unroll
+    for (int i = 0; i < 12; i++) {
+        st[i] = gf_add(st[i], sums[i & 3]);
+    }
+}
+
+// Internal linear layer: diagonal matrix plus full sum.
+__device__ __forceinline__ void int_layer(u64 st[12]) {
+    u64 sum = st[0];
+#pragma unroll
+    for (int i = 1; i < 12; i++) sum = gf_add(sum, st[i]);
+#pragma unroll
+    for (int i = 0; i < 12; i++) st[i] = gf_add(gf_mul(st[i], MDS_DIAG[i]), sum);
+}
+#endif
+
+#if LAIR_NOINLINE_PERMUTE
+// One copy of the (fully unrolled) permutation instead of three inlined ones:
+// 17k instructions instead of 51k, for instruction-cache pressure.
+__device__ __noinline__ void permute(u64 st[12]) {
+#else
+__device__ __forceinline__ void permute(u64 st[12]) {
+#endif
+#if LAIR_LAZY_ADD && LAIR_FOLD_RC
+    // Each linear layer folds the constant of the round that follows it.
+    ext_layer_rc(st, RC_INITIAL[0]);
+#pragma unroll
+    for (int r = 0; r < N_EXTERNAL_HALF; r++) {
+#pragma unroll
+        for (int i = 0; i < 12; i++) st[i] = gf_sbox(st[i]);
+        if (r + 1 < N_EXTERNAL_HALF) {
+            ext_layer_rc(st, RC_INITIAL[r + 1]);
+        } else {
+            ext_layer(st);
+            st[0] = gf_add(st[0], RC_INTERNAL[0]);
+        }
+    }
+    LAIR_UNROLL(LAIR_INT_UNROLL)
+    for (int r = 0; r < N_INTERNAL; r++) {
+        st[0] = gf_sbox(st[0]);
+        if (r + 1 < N_INTERNAL) {
+            int_layer_rc(st, RC_INTERNAL[r + 1], true, nullptr);
+        } else {
+            int_layer_rc(st, 0ull, false, RC_TERMINAL[0]);
+        }
+    }
+#pragma unroll
+    for (int r = 0; r < N_EXTERNAL_HALF; r++) {
+#pragma unroll
+        for (int i = 0; i < 12; i++) st[i] = gf_sbox(st[i]);
+        if (r + 1 < N_EXTERNAL_HALF) {
+            ext_layer_rc(st, RC_TERMINAL[r + 1]);
+        } else {
+            ext_layer(st);
+        }
+    }
+#else
+    ext_layer(st);
+#pragma unroll
+    for (int r = 0; r < N_EXTERNAL_HALF; r++) {
+#pragma unroll
+        for (int i = 0; i < 12; i++) st[i] = gf_sbox(gf_add(st[i], RC_INITIAL[r][i]));
+        ext_layer(st);
+    }
+    LAIR_UNROLL(LAIR_INT_UNROLL)
+    for (int r = 0; r < N_INTERNAL; r++) {
+        st[0] = gf_sbox(gf_add(st[0], RC_INTERNAL[r]));
+        int_layer(st);
+    }
+#pragma unroll
+    for (int r = 0; r < N_EXTERNAL_HALF; r++) {
+#pragma unroll
+        for (int i = 0; i < 12; i++) st[i] = gf_sbox(gf_add(st[i], RC_TERMINAL[r][i]));
+        ext_layer(st);
+    }
+#endif
+}
+
+__device__ __forceinline__ u32 bswap32(u32 v) {
+    return __byte_perm(v, 0, 0x0123);
+}
+
+// Bindings mirror the WGSL kernel:
+//   results  [flag, nonce(16 u32 LE), hash(16 u32 LE)]  = 33 u32
+//   midstate 12 felts as u64
+//   start_nonce 16 u32 LE limbs
+//   target 16 u32 LE limbs
+// lair: the launch-uniform inputs ride in the parameter bank instead of being
+// copied into per-thread registers (quantus/miner#3). Keeping midstate (24
+// registers), target (16) and the nonce base (16) live across the whole nonce
+// loop cost 56 registers of residency that the permutation's own working set
+// needs more: ptxas spilled 104 bytes per thread on sm_120 and 44 on sm_86/89.
+// Parameters are constant memory -- broadcast and cached, no register cost --
+// and they arrive with the launch, so the host no longer copies midstate,
+// start_nonce or target to the device once per batch.
+struct MiningUniforms {
+    u64 midstate[12];
+    u32 start_nonce[16];
+    u32 target[16];
+};
+static_assert(sizeof(MiningUniforms) == 224, "MiningUniforms must match the host layout");
+
+extern "C" __global__ void __launch_bounds__(256)
+mining_main(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;
+    if (tid >= total_threads) return;
+    u32 base_index = tid * nonces_per_thread;
+
+    for (u32 j = 0; j < nonces_per_thread; j++) {
+        u32 logical_index = base_index + j;
+        if (logical_index >= total_nonces) break;
+        if (j > 0u && *((volatile u32*)results) != 0u) return;
+
+        // Low 256 bits only; the host guarantees no carry into limbs 8..15, so
+        // the high half is never materialised here -- on a hit it is written
+        // straight from the parameter bank.
+        u32 current_nonce[8];
+        u32 val0 = uni.start_nonce[0];
+        u32 sum0 = val0 + logical_index;
+        current_nonce[0] = sum0;
+        u32 carry = sum0 < val0 ? 1u : 0u;
+#pragma unroll
+        for (int i = 1; i < 8; i++) {
+            u32 val = uni.start_nonce[i];
+            u32 sum = val + carry;
+            current_nonce[i] = sum;
+            carry = sum < val ? 1u : 0u;
+        }
+
+        // Resume the sponge from the midstate: absorb the low nonce half,
+        // pad, squeeze twice (3 permutations instead of 5; the second squeeze
+        // only for candidates).
+        u64 st[12];
+#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);
+        st[0] = gf_add(st[0], 1ull);
+        st[1] = gf_add(st[1], 1ull);
+        permute(st);
+
+        // First squeeze: most significant 256 bits of the hash.
+        u32 first[8];
+#pragma unroll
+        for (int i = 0; i < 4; i++) {
+            u64 c = gf_canon(st[i]);
+            first[2 * i] = (u32)(c & EPS);
+            first[2 * i + 1] = (u32)(c >> 32);
+        }
+        u32 cmp = 0u; // 0 equal so far, 1 above target, 2 below target
+#pragma unroll
+        for (int i = 0; i < 8; i++) {
+            u32 h = bswap32(first[i]);
+            u32 t = uni.target[15 - i];
+            if (h != t) { cmp = h > t ? 1u : 2u; break; }
+        }
+        if (cmp == 1u) continue;
+
+        u32 hash_le[16];
+#pragma unroll
+        for (int i = 0; i < 8; i++) hash_le[15 - i] = bswap32(first[i]);
+        permute(st);
+#pragma unroll
+        for (int i = 0; i < 4; i++) {
+            u64 c = gf_canon(st[i]);
+            hash_le[7 - 2 * i] = bswap32((u32)(c & EPS));
+            hash_le[6 - 2 * i] = bswap32((u32)(c >> 32));
+        }
+        bool below = (cmp == 2u);
+        if (!below) {
+#pragma unroll
+            for (int i = 0; i < 8; i++) {
+                u32 h = hash_le[7 - i];
+                u32 t = uni.target[7 - i];
+                if (h != t) { below = h < t; break; }
+            }
+        }
+
+        if (below) {
+            if (atomicExch(&results[0], 1u) == 0u) {
+#pragma unroll
+                for (int i = 0; i < 8; i++) results[1 + i] = current_nonce[i];
+#pragma unroll
+                for (int i = 8; i < 16; i++) results[1 + i] = uni.start_nonce[i];
+#pragma unroll
+                for (int i = 0; i < 16; i++) results[17 + i] = hash_le[i];
+                __threadfence();
+            }
+            return;
+        }
+    }
+}

crates/engine-cuda/src/lib.rs

@@ -0,0 +1,509 @@
+#![deny(rust_2018_idioms)]
+
+//! lair: native CUDA mining engine behind `MinerEngine` (quantus/miner#3).
+//!
+//! Same host contract and batch loop as `engine-gpu` (midstate per batch, no
+//! carry into the high nonce half, cancellation between batches, thread-local
+//! worker-to-device assignment), with the kernel in `kernels/mining.cu` built
+//! into a fat binary by `build.rs`. Nothing in `engine-gpu` is touched; the
+//! service picks this engine when the binary carries a kernel and a CUDA
+//! driver is present, and falls back to wgpu otherwise.
+
+use cudarc::driver::{
+    CudaContext, CudaFunction, CudaSlice, CudaStream, DeviceRepr, DriverError, LaunchConfig,
+    PushKernelArg,
+};
+use cudarc::nvrtc::Ptx;
+use engine_cpu::{CancelCheck, Candidate, EngineStatus, FoundOrigin, MinerEngine, Range};
+use pow_core::{format_hashrate, format_u512, JobContext};
+use primitive_types::U512;
+use std::cell::RefCell;
+use std::sync::atomic::{AtomicUsize, Ordering};
+use std::sync::Arc;
+use std::time::Instant;
+
+/// The fat binary produced by build.rs; empty when nvcc was not available.
+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";
+const THREADS_PER_BLOCK: u32 = 256;
+/// Threads per SM the grid is sized for (override: MINER_CUDA_THREADS_PER_SM).
+/// Sets how many nonces each thread loops over for a given batch; tuned with #2.
+// Measured on the 4090 (#3): 8192 caps the grid at ~1M threads, so batches
+// above 1M loop nonces per thread and lose up to 4%; 32768 keeps one nonce per
+// thread up to 4M and is flat across batch sizes at ~305 MH/s.
+const THREADS_PER_SM_DEFAULT: u32 = 32768;
+
+fn threads_per_sm() -> u32 {
+    std::env::var("MINER_CUDA_THREADS_PER_SM")
+        .ok()
+        .and_then(|v| v.parse().ok())
+        .filter(|v: &u32| *v > 0)
+        .unwrap_or(THREADS_PER_SM_DEFAULT)
+}
+
+static ENGINE_ID_COUNTER: AtomicUsize = AtomicUsize::new(0);
+
+thread_local! {
+    static ASSIGNED_DEVICE: RefCell<Option<(usize, usize)>> = const { RefCell::new(None) };
+    static WORKER_RESOURCES: RefCell<Option<(usize, WorkerResources)>> = const { RefCell::new(None) };
+    static DEVICE_LOST: RefCell<Option<usize>> = const { RefCell::new(None) };
+}
+
+struct Device {
+    ctx: Arc<CudaContext>,
+    func: CudaFunction,
+    name: String,
+    sm_count: u32,
+    threads_per_sm: u32,
+    metrics: metrics::DeviceMetrics,
+}
+
+struct WorkerResources {
+    stream: Arc<CudaStream>,
+    results: CudaSlice<u32>,
+    host_results: Vec<u32>,
+    /// Target limbs for this job context, recomputed when the target changes.
+    target_u32s: [u32; 16],
+    /// Target the limbs above were derived from.
+    target_written: Option<U512>,
+}
+
+const RESULTS_LEN: usize = 1 + 16 + 16;
+
+/// lair: the kernel's launch-uniform inputs, passed by value so they land in
+/// the parameter bank (constant memory) instead of per-thread registers.
+///
+/// Layout must match `struct MiningUniforms` in `kernels/mining.cu`. Passing
+/// them as parameters rather than device buffers took the kernel from 128
+/// registers with 104 bytes of per-thread spill on sm_120 (44 on sm_86/89) to
+/// 106 / 80 registers with none, and removed three host-to-device copies per
+/// batch.
+#[repr(C)]
+#[derive(Clone, Copy)]
+struct MiningUniforms {
+    midstate: [u64; 12],
+    start_nonce: [u32; 16],
+    target: [u32; 16],
+}
+
+// SAFETY: `#[repr(C)]` plain-old-data with no padding (12 x u64 then 32 x u32),
+// 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 {}
+
+// 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>() == 224);
+
+pub struct CudaEngine {
+    engine_id: usize,
+    devices: Vec<Arc<Device>>,
+    device_counter: AtomicUsize,
+    batch_size: u32,
+    throttle_ms: u64,
+}
+
+impl CudaEngine {
+    /// Returns an error when the binary carries no kernel, no CUDA driver is
+    /// present, or no device initialises.
+    pub fn try_new(batch_size: u32, throttle_ms: u64) -> Result<Self, String> {
+        if batch_size == 0 {
+            return Err("batch size must be non-zero".into());
+        }
+        if FATBIN.is_empty() {
+            return Err("CUDA kernel not compiled into this binary (built without nvcc)".into());
+        }
+        let count =
+            CudaContext::device_count().map_err(|e| format!("CUDA driver unavailable: {e:?}"))?;
+        if count <= 0 {
+            return Err("no CUDA devices".into());
+        }
+        let mut devices = Vec::new();
+        for ordinal in 0..count as usize {
+            match Device::init(ordinal) {
+                Ok(d) => {
+                    log::info!(
+                        target: "cuda_engine",
+                        "CUDA device {ordinal}: {} ({} SMs) using {KERNEL_ID} kernel [sm_{{{}}}]",
+                        d.name,
+                        d.sm_count,
+                        CUDA_ARCHS
+                    );
+                    devices.push(Arc::new(d));
+                }
+                Err(e) => {
+                    log::warn!(target: "cuda_engine", "CUDA device {ordinal} failed to initialise: {e:?}; skipping");
+                }
+            }
+        }
+        if devices.is_empty() {
+            return Err("no CUDA device could be initialised".into());
+        }
+        log::info!(
+            target: "cuda_engine",
+            "CUDA engine initialized with {} devices (batch size: {} nonces, throttle: {}ms)",
+            devices.len(),
+            batch_size,
+            throttle_ms
+        );
+        Ok(Self {
+            engine_id: ENGINE_ID_COUNTER.fetch_add(1, Ordering::SeqCst),
+            devices,
+            device_counter: AtomicUsize::new(0),
+            batch_size,
+            throttle_ms,
+        })
+    }
+
+    pub fn device_count(&self) -> usize {
+        self.devices.len()
+    }
+
+    /// Drop the calling thread's device resources (call on worker exit).
+    pub fn clear_worker_resources() {
+        WORKER_RESOURCES.with(|r| *r.borrow_mut() = None);
+        ASSIGNED_DEVICE.with(|a| *a.borrow_mut() = None);
+    }
+}
+
+impl Device {
+    fn init(ordinal: usize) -> Result<Self, DriverError> {
+        let ctx = CudaContext::new(ordinal)?;
+        let name = ctx.name()?;
+        let sm_count = ctx.attribute(
+            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")?;
+        Ok(Self {
+            ctx,
+            func,
+            name,
+            sm_count,
+            threads_per_sm: threads_per_sm(),
+            metrics: metrics::DeviceMetrics::new(ordinal, KERNEL_ID),
+        })
+    }
+
+    fn create_resources(&self) -> Result<WorkerResources, DriverError> {
+        let stream = self.ctx.new_stream()?;
+        Ok(WorkerResources {
+            results: stream.alloc_zeros::<u32>(RESULTS_LEN)?,
+            host_results: vec![0u32; RESULTS_LEN],
+            target_u32s: [0u32; 16],
+            target_written: None,
+            stream,
+        })
+    }
+}
+
+enum BatchResult {
+    Found {
+        candidate: Candidate,
+        hash_count: u64,
+    },
+    NotFound {
+        hash_count: u64,
+    },
+    DeviceLost,
+}
+
+impl MinerEngine for CudaEngine {
+    fn name(&self) -> &'static str {
+        "gpu-cuda"
+    }
+
+    fn prepare_context(&self, header_hash: [u8; 32], difficulty: U512) -> JobContext {
+        JobContext::new(header_hash, difficulty)
+    }
+
+    fn as_any(&self) -> &dyn std::any::Any {
+        self
+    }
+
+    fn search_range(
+        &self,
+        ctx: &JobContext,
+        range: Range,
+        cancel: &dyn CancelCheck,
+    ) -> EngineStatus {
+        if DEVICE_LOST.with(|l| *l.borrow() == Some(self.engine_id)) {
+            return EngineStatus::DeviceLost { hash_count: 0 };
+        }
+        if range.start > range.end {
+            return EngineStatus::Exhausted { hash_count: 0 };
+        }
+        if cancel.is_cancelled() {
+            return EngineStatus::Cancelled { hash_count: 0 };
+        }
+
+        let device_index = ASSIGNED_DEVICE.with(|a| {
+            let mut a = a.borrow_mut();
+            match *a {
+                Some((id, idx)) if id == self.engine_id => idx,
+                _ => {
+                    let idx = if self.devices.len() == 1 {
+                        0
+                    } else {
+                        self.device_counter.fetch_add(1, Ordering::SeqCst) % self.devices.len()
+                    };
+                    *a = Some((self.engine_id, idx));
+                    log::info!(target: "cuda_engine", "Worker thread assigned to CUDA device {idx} (of {} total devices)", self.devices.len());
+                    idx
+                }
+            }
+        });
+        let dev = &self.devices[device_index];
+
+        WORKER_RESOURCES.with(|cell| {
+            let mut slot = cell.borrow_mut();
+            let need_new = !matches!(&*slot, Some((id, _)) if *id == self.engine_id);
+            if need_new {
+                match dev.create_resources() {
+                    Ok(r) => *slot = Some((self.engine_id, r)),
+                    Err(e) => {
+                        log::error!(target: "cuda_engine", "CUDA device {device_index} resource allocation failed: {e:?}");
+                        DEVICE_LOST.with(|l| *l.borrow_mut() = Some(self.engine_id));
+                        return EngineStatus::DeviceLost { hash_count: 0 };
+                    }
+                }
+            }
+            let (_, res) = slot.as_mut().expect("resources present");
+            self.search_on(dev, device_index, res, ctx, range, cancel)
+        })
+    }
+}
+
+impl CudaEngine {
+    fn search_on(
+        &self,
+        dev: &Device,
+        device_index: usize,
+        res: &mut WorkerResources,
+        ctx: &JobContext,
+        range: Range,
+        cancel: &dyn CancelCheck,
+    ) -> EngineStatus {
+        // Target is per job; convert it to limbs once per context. It reaches
+        // the kernel in the parameter bank, so there is nothing to upload.
+        if res.target_written != Some(ctx.target) {
+            let target_bytes = ctx.target.to_little_endian();
+            for i in 0..16 {
+                res.target_u32s[i] =
+                    u32::from_le_bytes(target_bytes[i * 4..(i + 1) * 4].try_into().unwrap());
+            }
+            res.target_written = Some(ctx.target);
+        }
+
+        let search_start = Instant::now();
+        let mut total_hashes: u64 = 0;
+        let mut current_start = range.start;
+        let mut batch_num = 0u64;
+        let mut last_batch_hashes: u64 = 0;
+
+        log::info!(
+            target: "cuda_engine",
+            "CUDA {} search started: range {}..{}, batch size: {} nonces",
+            device_index,
+            format_u512(range.start),
+            format_u512(range.end),
+            self.batch_size
+        );
+
+        while current_start <= range.end {
+            if cancel.is_cancelled() {
+                dev.metrics.record_stale_hashes(last_batch_hashes);
+                let elapsed = search_start.elapsed();
+                log::info!(
+                    target: "cuda_engine",
+                    "CUDA {} cancelled before batch {} ({} total hashes in {:.2}s, {})",
+                    device_index,
+                    batch_num,
+                    total_hashes,
+                    elapsed.as_secs_f64(),
+                    format_hashrate(total_hashes as f64 / elapsed.as_secs_f64())
+                );
+                return EngineStatus::Cancelled {
+                    hash_count: total_hashes,
+                };
+            }
+
+            // Clamp so nonce increments never carry into the high 256 bits,
+            // which the midstate precompute relies on.
+            let remaining = range
+                .end
+                .saturating_sub(current_start)
+                .saturating_add(U512::one());
+            let headroom =
+                (U512::one() << 256) - (current_start & ((U512::one() << 256) - U512::one()));
+            let cap = remaining.min(headroom);
+            let this_batch: u32 = if cap > U512::from(self.batch_size) {
+                self.batch_size
+            } else {
+                cap.low_u32()
+            };
+
+            match self.run_single_batch(dev, res, ctx, current_start, this_batch) {
+                BatchResult::Found {
+                    candidate,
+                    hash_count,
+                } => {
+                    total_hashes += hash_count;
+                    return EngineStatus::Found {
+                        candidate,
+                        hash_count: total_hashes,
+                        origin: FoundOrigin::GpuG1,
+                    };
+                }
+                BatchResult::NotFound { hash_count } => {
+                    total_hashes += hash_count;
+                    last_batch_hashes = hash_count;
+                }
+                BatchResult::DeviceLost => return self.device_lost(dev, total_hashes),
+            }
+
+            current_start = current_start.saturating_add(U512::from(this_batch));
+            batch_num += 1;
+
+            if self.throttle_ms > 0 && current_start <= range.end {
+                let step = std::time::Duration::from_millis((self.throttle_ms / 10).max(1));
+                let mut remaining = std::time::Duration::from_millis(self.throttle_ms);
+                while remaining > std::time::Duration::ZERO {
+                    if cancel.is_cancelled() {
+                        return EngineStatus::Cancelled {
+                            hash_count: total_hashes,
+                        };
+                    }
+                    let s = remaining.min(step);
+                    std::thread::sleep(s);
+                    remaining = remaining.saturating_sub(s);
+                }
+            }
+        }
+
+        let elapsed = search_start.elapsed();
+        log::info!(
+            target: "cuda_engine",
+            "CUDA {} search exhausted: {} hashes in {} batches ({:.2}s, {})",
+            device_index,
+            total_hashes,
+            batch_num,
+            elapsed.as_secs_f64(),
+            format_hashrate(total_hashes as f64 / elapsed.as_secs_f64())
+        );
+        EngineStatus::Exhausted {
+            hash_count: total_hashes,
+        }
+    }
+
+    fn device_lost(&self, dev: &Device, hash_count: u64) -> EngineStatus {
+        dev.metrics.record_device_lost();
+        DEVICE_LOST.with(|l| *l.borrow_mut() = Some(self.engine_id));
+        log::error!(target: "cuda_engine", "CUDA device lost or unresponsive - stopping worker");
+        EngineStatus::DeviceLost { hash_count }
+    }
+
+    fn run_single_batch(
+        &self,
+        dev: &Device,
+        res: &mut WorkerResources,
+        ctx: &JobContext,
+        batch_start: U512,
+        batch_size: u32,
+    ) -> BatchResult {
+        let batch_start_at = Instant::now();
+
+        // Grid: enough threads to fill the card several waves deep, then loop
+        // the remainder per thread.
+        let max_threads = (dev.sm_count * dev.threads_per_sm) as u64;
+        let logical_threads = (batch_size as u64).min(max_threads).max(1);
+        let num_blocks = ((logical_threads as u32).div_ceil(THREADS_PER_BLOCK)).max(1);
+        let total_threads = (num_blocks * THREADS_PER_BLOCK) as u64;
+        let nonces_per_thread = ((batch_size as u64).div_ceil(total_threads)).max(1) as u32;
+        let total_threads_u32 = total_threads as u32;
+
+        let start_nonce_bytes = batch_start.to_little_endian();
+        let mut start_u32s = [0u32; 16];
+        for i in 0..16 {
+            start_u32s[i] =
+                u32::from_le_bytes(start_nonce_bytes[i * 4..(i + 1) * 4].try_into().unwrap());
+        }
+        let nonce_be = batch_start.to_big_endian();
+        let midstate = pow_core::mining_midstate(ctx.header, nonce_be[..32].try_into().unwrap());
+
+        let uni = MiningUniforms {
+            midstate,
+            start_nonce: start_u32s,
+            target: res.target_u32s,
+        };
+
+        let stream = res.stream.clone();
+        let r: Result<(), DriverError> = (|| {
+            stream.memset_zeros(&mut res.results)?;
+            let cfg = LaunchConfig {
+                grid_dim: (num_blocks, 1, 1),
+                block_dim: (THREADS_PER_BLOCK, 1, 1),
+                shared_mem_bytes: 0,
+            };
+            let mut launch = stream.launch_builder(&dev.func);
+            launch
+                .arg(&mut res.results)
+                .arg(&uni)
+                .arg(&total_threads_u32)
+                .arg(&nonces_per_thread)
+                .arg(&batch_size);
+            // SAFETY: the kernel signature matches the argument list above and
+            // every pointer is a device buffer of at least the size the kernel reads.
+            unsafe { launch.launch(cfg) }?;
+            Ok(())
+        })();
+        if let Err(e) = r {
+            log::error!(target: "cuda_engine", "CUDA batch submit failed: {e:?}");
+            return BatchResult::DeviceLost;
+        }
+        let submitted_at = Instant::now();
+
+        if let Err(e) = stream
+            .memcpy_dtoh(&res.results, &mut res.host_results)
+            .and_then(|_| stream.synchronize())
+        {
+            log::error!(target: "cuda_engine", "CUDA batch readback failed: {e:?}");
+            return BatchResult::DeviceLost;
+        }
+        let gpu_time = submitted_at.elapsed();
+        dev.metrics
+            .observe_batch(gpu_time, batch_start_at.elapsed() - gpu_time);
+
+        let out = &res.host_results;
+        let dispatched = (total_threads * nonces_per_thread as u64).min(batch_size as u64);
+        if out[0] != 0 {
+            let nonce = U512::from_little_endian(&bytemuck_cast(&out[1..17]));
+            let hash = U512::from_little_endian(&bytemuck_cast(&out[17..33]));
+            let work = nonce.to_big_endian();
+            let hashes_computed = if nonce >= batch_start {
+                let logical_index = (nonce - batch_start).as_u64();
+                let winning_iteration = logical_index % (nonces_per_thread as u64);
+                (total_threads * (winning_iteration + 1)).min(dispatched)
+            } else {
+                dispatched
+            };
+            dev.metrics.record_hashes(hashes_computed);
+            dev.metrics.record_solution();
+            return BatchResult::Found {
+                candidate: Candidate { nonce, work, hash },
+                hash_count: hashes_computed,
+            };
+        }
+        dev.metrics.record_hashes(dispatched);
+        BatchResult::NotFound {
+            hash_count: dispatched,
+        }
+    }
+}
+
+fn bytemuck_cast(words: &[u32]) -> Vec<u8> {
+    words.iter().flat_map(|w| w.to_le_bytes()).collect()
+}

crates/engine-gpu/Cargo.toml

@@ -19,6 +19,7 @@ simd-poseidon2 = []
 
 [dependencies]
 engine-cpu = { path = "../engine-cpu" }
+metrics = { path = "../metrics" } # lair: per-device metrics (quantus/miner#9)
 pow-core = { path = "../pow-core" }
 primitive-types = { workspace = true }
 log = { workspace = true }

crates/engine-gpu/src/lib.rs

@@ -20,6 +20,8 @@ use std::sync::{
 
 /// Represents a single GPU device context.
 struct GpuContext {
+    // lair: per-device metric handles, labelled with the kernel id (quantus/miner#9).
+    metrics: metrics::DeviceMetrics,
     device: wgpu::Device,
     queue: wgpu::Queue,
     pipeline: wgpu::ComputePipeline,
@@ -443,6 +445,7 @@ impl GpuEngine {
 
             initialized.push(InitializedGpu {
                 context: Arc::new(GpuContext {
+                    metrics: metrics::DeviceMetrics::new(i, kernel.id()),
                     device,
                     queue,
                     pipeline,
@@ -615,6 +618,9 @@ impl MinerEngine for GpuEngine {
         let mut total_hashes: u64 = 0;
         let mut current_start = range.start;
         let mut batch_num = 0u64;
+        // lair: hashes of the most recent batch; if the job turns out to have
+        // been superseded while it ran, that batch was wasted work.
+        let mut last_batch_hashes: u64 = 0;
 
         log::info!(
             target: "gpu_engine",
@@ -629,6 +635,7 @@ impl MinerEngine for GpuEngine {
         while current_start <= range.end {
             // Check for cancellation at host level BEFORE starting each batch
             if cancel.is_cancelled() {
+                gpu_ctx.metrics.record_stale_hashes(last_batch_hashes); // lair
                 let elapsed = search_start.elapsed();
                 let hash_rate = total_hashes as f64 / elapsed.as_secs_f64();
                 log::info!(
@@ -696,10 +703,12 @@ impl MinerEngine for GpuEngine {
                 }
                 BatchResult::NotFound { hash_count } => {
                     total_hashes += hash_count;
+                    last_batch_hashes = hash_count; // lair
                 }
                 BatchResult::DeviceLost => {
-                    // GPU device is lost/unresponsive - mark as permanently dead
-                    // and clear resources to prevent "buffer already mapped" panics
+                    gpu_ctx.metrics.record_device_lost(); // lair
+                                                          // GPU device is lost/unresponsive - mark as permanently dead
+                                                          // and clear resources to prevent "buffer already mapped" panics
                     DEVICE_LOST.with(|lost| *lost.borrow_mut() = Some(self.engine_id));
                     WORKER_RESOURCES.with(|res| *res.borrow_mut() = None);
 
@@ -793,7 +802,8 @@ fn run_single_batch(
     batch_start: U512,
     batch_size: u32,
 ) -> BatchResult {
-    // Calculate dispatch configuration for this batch
+    let batch_start_at = std::time::Instant::now(); // lair: batch phase timing
+                                                    // Calculate dispatch configuration for this batch
     let threads_per_workgroup = 256u32;
     let limits = gpu_ctx.device.limits();
     let max_workgroups = limits.max_compute_workgroups_per_dimension;
@@ -865,6 +875,7 @@ fn run_single_batch(
     );
 
     gpu_ctx.queue.submit(Some(encoder.finish()));
+    let submitted_at = std::time::Instant::now(); // lair
 
     // Wait for GPU to complete (blocking)
     let buffer_slice = resources.staging_buffer.slice(..);
@@ -911,6 +922,12 @@ fn run_single_batch(
 
     // Only reach here if final_status == 1 (success), buffer is mapped
     debug_assert_eq!(final_status, 1);
+    // lair: gpu = submit to mapped; host = the rest of the batch so far plus
+    // the readback below, which is small and constant.
+    let gpu_time = submitted_at.elapsed();
+    gpu_ctx
+        .metrics
+        .observe_batch(gpu_time, batch_start_at.elapsed() - gpu_time);
 
     // Read results
     let data = buffer_slice.get_mapped_range();
@@ -939,6 +956,8 @@ fn run_single_batch(
         drop(data);
         resources.staging_buffer.unmap();
 
+        gpu_ctx.metrics.record_hashes(hashes_computed); // lair
+        gpu_ctx.metrics.record_solution(); // lair
         return BatchResult::Found {
             candidate: Candidate { nonce, work, hash },
             hash_count: hashes_computed,
@@ -948,6 +967,7 @@ fn run_single_batch(
     drop(data);
     resources.staging_buffer.unmap();
 
+    gpu_ctx.metrics.record_hashes(dispatched_nonces); // lair
     BatchResult::NotFound {
         hash_count: dispatched_nonces,
     }

crates/metrics/src/lair.rs

@@ -0,0 +1,344 @@
+//! lair: metrics that attribute performance to builds, devices and jobs
+//! (quantus/miner#9). Additive to origin's metrics in `lib.rs`; nothing there
+//! changes, so the fleet dashboard keeps working across origin merges.
+//!
+//! Naming: `miner_*` like origin's. Per-device series carry `device` (engine
+//! index) and `kernel` (the kernel id the engine selected), so a silent
+//! fallback from one kernel to another shows up as a label change rather than
+//! as an unexplained hashrate drop.
+
+use crate::REGISTRY;
+use once_cell::sync::Lazy;
+use prometheus::{
+    Counter, Histogram, HistogramOpts, HistogramVec, IntCounter, IntCounterVec, IntGauge,
+    IntGaugeVec, Opts,
+};
+use std::time::Duration;
+
+fn reg<M: prometheus::core::Collector + Clone + 'static>(m: M) -> M {
+    REGISTRY
+        .register(Box::new(m.clone()))
+        .expect("register lair metric");
+    m
+}
+
+// ---------------------------------------------------------------------------
+// Identity: which build and which configuration produced the numbers
+// ---------------------------------------------------------------------------
+
+static BUILD_INFO: Lazy<IntGaugeVec> = Lazy::new(|| {
+    reg(IntGaugeVec::new(
+        Opts::new(
+            "miner_build_info",
+            "Build identity of the running miner, always 1",
+        ),
+        &["version", "commit"],
+    )
+    .expect("miner_build_info"))
+});
+
+static CONFIG_INFO: Lazy<IntGaugeVec> = Lazy::new(|| {
+    reg(IntGaugeVec::new(
+        Opts::new(
+            "miner_config_info",
+            "Effective mining configuration, always 1",
+        ),
+        &[
+            "engine",
+            "gpu_batch_size",
+            "gpu_devices",
+            "cpu_workers",
+            "gpu_throttle_ms",
+        ],
+    )
+    .expect("miner_config_info"))
+});
+
+/// Set once at startup. `commit` is the SHA embedded by build.rs.
+pub fn set_build_info(version: &str, commit: &str) {
+    BUILD_INFO.with_label_values(&[version, commit]).set(1);
+}
+
+/// Set once the engines are resolved. Two deploys of one commit with different
+/// flags must not look like the same experiment.
+pub fn set_config_info(
+    engine: &str,
+    gpu_batch_size: u32,
+    gpu_devices: usize,
+    cpu_workers: usize,
+    gpu_throttle_ms: u64,
+) {
+    CONFIG_INFO
+        .with_label_values(&[
+            engine,
+            &gpu_batch_size.to_string(),
+            &gpu_devices.to_string(),
+            &cpu_workers.to_string(),
+            &gpu_throttle_ms.to_string(),
+        ])
+        .set(1);
+}
+
+// ---------------------------------------------------------------------------
+// Per device: throughput, outcomes, and where the time goes inside a batch
+// ---------------------------------------------------------------------------
+
+static DEVICE_HASHES: Lazy<IntCounterVec> = Lazy::new(|| {
+    reg(IntCounterVec::new(
+        Opts::new(
+            "miner_device_hashes_total",
+            "Hashes computed, per GPU device and kernel",
+        ),
+        &["device", "kernel"],
+    )
+    .expect("miner_device_hashes_total"))
+});
+
+static DEVICE_SOLUTIONS: Lazy<IntCounterVec> = Lazy::new(|| {
+    reg(IntCounterVec::new(
+        Opts::new(
+            "miner_device_solutions_total",
+            "Solutions found, per GPU device and kernel",
+        ),
+        &["device", "kernel"],
+    )
+    .expect("miner_device_solutions_total"))
+});
+
+static DEVICE_LOST: Lazy<IntCounterVec> = Lazy::new(|| {
+    reg(IntCounterVec::new(
+        Opts::new(
+            "miner_device_lost_total",
+            "Times a GPU device was lost or unresponsive and its worker stopped",
+        ),
+        &["device", "kernel"],
+    )
+    .expect("miner_device_lost_total"))
+});
+
+static DEVICE_STALE_HASHES: Lazy<IntCounterVec> = Lazy::new(|| {
+    reg(IntCounterVec::new(
+        Opts::new(
+            "miner_stale_hashes_total",
+            "Hashes computed after the job they were for had been superseded: the batch in \
+             flight when a new job arrived. The wasted-work cost of batch size.",
+        ),
+        &["device", "kernel"],
+    )
+    .expect("miner_stale_hashes_total"))
+});
+
+static GPU_BATCH_SECONDS: Lazy<HistogramVec> = Lazy::new(|| {
+    reg(HistogramVec::new(
+        HistogramOpts::new(
+            "miner_gpu_batch_seconds",
+            "Per-batch time split into the GPU executing (phase=gpu) and the host \
+             preparing, submitting and reading back (phase=host)",
+        )
+        .buckets(vec![
+            0.0005, 0.001, 0.002, 0.004, 0.008, 0.016, 0.032, 0.064, 0.128, 0.256, 0.512, 1.0, 2.0,
+        ]),
+        &["device", "kernel", "phase"],
+    )
+    .expect("miner_gpu_batch_seconds"))
+});
+
+/// Handles for one device, resolved once so the per-batch path does no label
+/// lookups. Created by the engine when it initialises a device.
+#[derive(Clone)]
+pub struct DeviceMetrics {
+    hashes: IntCounter,
+    solutions: IntCounter,
+    lost: IntCounter,
+    stale: IntCounter,
+    batch_gpu: Histogram,
+    batch_host: Histogram,
+}
+
+impl DeviceMetrics {
+    pub fn new(device: usize, kernel: &str) -> Self {
+        let d = device.to_string();
+        Self {
+            hashes: DEVICE_HASHES.with_label_values(&[&d, kernel]),
+            solutions: DEVICE_SOLUTIONS.with_label_values(&[&d, kernel]),
+            lost: DEVICE_LOST.with_label_values(&[&d, kernel]),
+            stale: DEVICE_STALE_HASHES.with_label_values(&[&d, kernel]),
+            batch_gpu: GPU_BATCH_SECONDS.with_label_values(&[&d, kernel, "gpu"]),
+            batch_host: GPU_BATCH_SECONDS.with_label_values(&[&d, kernel, "host"]),
+        }
+    }
+
+    pub fn record_hashes(&self, n: u64) {
+        self.hashes.inc_by(n);
+    }
+
+    pub fn record_solution(&self) {
+        self.solutions.inc();
+    }
+
+    pub fn record_device_lost(&self) {
+        self.lost.inc();
+    }
+
+    /// The batch that completed after its job was superseded.
+    pub fn record_stale_hashes(&self, n: u64) {
+        self.stale.inc_by(n);
+    }
+
+    /// `gpu` is submit-to-completion on the device; `host` is everything else
+    /// in the batch (buffer writes, encoding, readback, bookkeeping).
+    pub fn observe_batch(&self, gpu: Duration, host: Duration) {
+        self.batch_gpu.observe(gpu.as_secs_f64());
+        self.batch_host.observe(host.as_secs_f64());
+    }
+}
+
+// ---------------------------------------------------------------------------
+// Jobs and results: the efficiency side of every throughput trade-off
+// ---------------------------------------------------------------------------
+
+static JOBS_RECEIVED: Lazy<IntCounter> = Lazy::new(|| {
+    reg(
+        IntCounter::new("miner_jobs_received_total", "Jobs received from the node")
+            .expect("miner_jobs_received_total"),
+    )
+});
+
+static JOB_PICKUP_SECONDS: Lazy<HistogramVec> = Lazy::new(|| {
+    reg(HistogramVec::new(
+        HistogramOpts::new(
+            "miner_job_pickup_seconds",
+            "Job issued to a worker starting on it. For a busy worker this is the \
+             time to notice cancellation and finish the in-flight batch",
+        )
+        .buckets(vec![
+            0.001, 0.002, 0.005, 0.01, 0.02, 0.05, 0.1, 0.2, 0.5, 1.0, 2.0, 5.0,
+        ]),
+        &["engine"],
+    )
+    .expect("miner_job_pickup_seconds"))
+});
+
+static RESULTS_SUBMITTED: Lazy<IntCounter> = Lazy::new(|| {
+    reg(IntCounter::new(
+        "miner_results_submitted_total",
+        "Solutions sent to the node",
+    )
+    .expect("miner_results_submitted_total"))
+});
+
+static RESULTS_SEND_FAILED: Lazy<IntCounter> = Lazy::new(|| {
+    reg(IntCounter::new(
+        "miner_results_send_failed_total",
+        "Solutions that could not be sent to the node",
+    )
+    .expect("miner_results_send_failed_total"))
+});
+
+static SEAL_LATENCY: Lazy<Histogram> = Lazy::new(|| {
+    reg(Histogram::with_opts(
+        HistogramOpts::new(
+            "miner_seal_latency_seconds",
+            "Solution found on a worker to the result sent to the node",
+        )
+        .buckets(vec![
+            0.001, 0.002, 0.005, 0.01, 0.02, 0.05, 0.1, 0.2, 0.5, 1.0, 2.0, 5.0,
+        ]),
+    )
+    .expect("miner_seal_latency_seconds"))
+});
+
+static JOB_IDLE_SECONDS: Lazy<Counter> = Lazy::new(|| {
+    reg(Counter::new(
+        "miner_job_idle_seconds_total",
+        "Seconds spent between sending a result and receiving the next job (node-attributable idle)",
+    )
+    .expect("miner_job_idle_seconds_total"))
+});
+
+pub fn record_job_received() {
+    JOBS_RECEIVED.inc();
+}
+
+pub fn observe_job_pickup(engine: &str, d: Duration) {
+    JOB_PICKUP_SECONDS
+        .with_label_values(&[engine])
+        .observe(d.as_secs_f64());
+}
+
+pub fn record_result_submitted(found_to_sent: Option<Duration>) {
+    RESULTS_SUBMITTED.inc();
+    if let Some(d) = found_to_sent {
+        SEAL_LATENCY.observe(d.as_secs_f64());
+    }
+}
+
+pub fn record_result_send_failed() {
+    RESULTS_SEND_FAILED.inc();
+}
+
+pub fn record_job_idle(d: Duration) {
+    JOB_IDLE_SECONDS.inc_by(d.as_secs_f64());
+}
+
+// ---------------------------------------------------------------------------
+// Connection to the node: is a hashrate drop the miner or the node?
+// ---------------------------------------------------------------------------
+
+static CONNECTS: Lazy<IntCounter> = Lazy::new(|| {
+    reg(IntCounter::new(
+        "miner_connects_total",
+        "Connections established to the node",
+    )
+    .expect("miner_connects_total"))
+});
+
+static CONNECT_FAILURES: Lazy<IntCounter> = Lazy::new(|| {
+    reg(IntCounter::new(
+        "miner_connect_failures_total",
+        "Connection attempts to the node that failed",
+    )
+    .expect("miner_connect_failures_total"))
+});
+
+static DISCONNECTS: Lazy<IntCounter> = Lazy::new(|| {
+    reg(IntCounter::new(
+        "miner_disconnects_total",
+        "Connections to the node that were lost",
+    )
+    .expect("miner_disconnects_total"))
+});
+
+static CONNECTED: Lazy<IntGauge> = Lazy::new(|| {
+    reg(
+        IntGauge::new("miner_connected", "1 while connected to the node, else 0")
+            .expect("miner_connected"),
+    )
+});
+
+static DISCONNECTED_SECONDS: Lazy<Counter> = Lazy::new(|| {
+    reg(Counter::new(
+        "miner_disconnected_seconds_total",
+        "Seconds spent without a node connection after the first successful connection",
+    )
+    .expect("miner_disconnected_seconds_total"))
+});
+
+pub fn record_connected() {
+    CONNECTS.inc();
+    CONNECTED.set(1);
+}
+
+pub fn record_connect_failed() {
+    CONNECT_FAILURES.inc();
+    CONNECTED.set(0);
+}
+
+pub fn record_disconnected() {
+    DISCONNECTS.inc();
+    CONNECTED.set(0);
+}
+
+pub fn record_disconnected_time(d: Duration) {
+    DISCONNECTED_SECONDS.inc_by(d.as_secs_f64());
+}

crates/metrics/src/lib.rs

@@ -39,6 +39,10 @@ use anyhow::Result;
 
 static REGISTRY: Lazy<Registry> = Lazy::new(Registry::new);
 
+// lair: build, device and job metrics (quantus/miner#9); see lair.rs.
+mod lair;
+pub use lair::*;
+
 // ---------------------------------------------------------------------------
 // Hash Rate Metrics
 // ---------------------------------------------------------------------------

crates/miner-cli/build.rs

@@ -0,0 +1,51 @@
+// lair: embed the git commit in the binary so `--version` and the build-info
+// metric identify a deployed build by commit, not by the workspace semver
+// (which does not change between commits on a branch that deploys on push).
+//
+// Resolution order:
+//   1. MINER_BUILD_SHA in the environment (CI sets it from the checked-out ref)
+//   2. `git rev-parse HEAD` of the workspace, with "-dirty" if the tree differs
+//   3. "unknown"
+use std::process::Command;
+
+fn git(args: &[&str]) -> Option<String> {
+    let out = Command::new("git").args(args).output().ok()?;
+    if !out.status.success() {
+        return None;
+    }
+    let s = String::from_utf8(out.stdout).ok()?;
+    let s = s.trim();
+    if s.is_empty() {
+        None
+    } else {
+        Some(s.to_string())
+    }
+}
+
+fn main() {
+    println!("cargo:rerun-if-env-changed=MINER_BUILD_SHA");
+
+    let sha = match std::env::var("MINER_BUILD_SHA") {
+        Ok(s) if !s.trim().is_empty() => s.trim().to_string(),
+        _ => match git(&["rev-parse", "--short=12", "HEAD"]) {
+            Some(head) => {
+                // Re-run when HEAD moves so a rebuild after a commit picks it up.
+                if let Some(dir) = git(&["rev-parse", "--git-dir"]) {
+                    println!("cargo:rerun-if-changed={dir}/HEAD");
+                    println!("cargo:rerun-if-changed={dir}/refs/heads");
+                }
+                let dirty = git(&["status", "--porcelain", "--untracked-files=no"])
+                    .map(|s| !s.is_empty())
+                    .unwrap_or(false);
+                if dirty {
+                    format!("{head}-dirty")
+                } else {
+                    head
+                }
+            }
+            None => "unknown".to_string(),
+        },
+    };
+
+    println!("cargo:rustc-env=MINER_BUILD_SHA={sha}");
+}

crates/miner-cli/src/main.rs

@@ -87,6 +87,11 @@ enum Command {
         #[arg(long = "allow-integrated", env = "MINER_ALLOW_INTEGRATED")]
         allow_integrated: bool,
 
+        /// lair: GPU engine: auto (CUDA when this binary carries a kernel and a
+        /// driver is present, else wgpu), cuda, or wgpu (quantus/miner#3).
+        #[arg(long = "gpu-engine", env = "MINER_GPU_ENGINE", default_value = "auto")]
+        gpu_engine: String,
+
         /// Enable verbose logging
         #[arg(short, long, env = "MINER_VERBOSE")]
         verbose: bool,
@@ -118,15 +123,29 @@ enum Command {
         #[arg(long = "allow-integrated", env = "MINER_ALLOW_INTEGRATED")]
         allow_integrated: bool,
 
+        /// lair: GPU engine: auto (CUDA when this binary carries a kernel and a
+        /// driver is present, else wgpu), cuda, or wgpu (quantus/miner#3).
+        #[arg(long = "gpu-engine", env = "MINER_GPU_ENGINE", default_value = "auto")]
+        gpu_engine: String,
+
         /// Enable verbose logging
         #[arg(short, long, env = "MINER_VERBOSE")]
         verbose: bool,
     },
 }
 
+// lair: semver plus the commit the binary was built from (see build.rs), so a
+// deploy can assert the running binary is the commit it shipped.
+const VERSION: &str = concat!(
+    env!("CARGO_PKG_VERSION"),
+    " (",
+    env!("MINER_BUILD_SHA"),
+    ")"
+);
+
 /// Quantus External Miner CLI
 #[derive(Parser, Debug)]
-#[command(author, version, about, long_about = None)]
+#[command(author, version = VERSION, about, long_about = None)]
 struct Args {
     #[command(subcommand)]
     command: Option<Command>,
@@ -135,6 +154,8 @@ struct Args {
 #[tokio::main]
 async fn main() {
     let args = Args::parse();
+    // lair: identify this build in the metrics (quantus/miner#9).
+    metrics::set_build_info(env!("CARGO_PKG_VERSION"), env!("MINER_BUILD_SHA"));
 
     let Some(command) = args.command else {
         eprintln!("Error: No command provided. Use 'serve' to start mining (defaults to local node at 127.0.0.1:9833).");
@@ -160,6 +181,7 @@ async fn main() {
             gpu_throttle_ms,
             metrics_port,
             allow_integrated,
+            gpu_engine,
             verbose,
         } => {
             init_logger(verbose);
@@ -207,6 +229,7 @@ async fn main() {
                 cpu_batch_size,
                 gpu_throttle_ms,
                 allow_integrated,
+                gpu_engine,
             };
 
             if let Err(e) = run(config).await {
@@ -222,6 +245,7 @@ async fn main() {
             cpu_batch_size,
             duration,
             allow_integrated,
+            gpu_engine,
             verbose,
         } => {
             init_logger(verbose);
@@ -232,6 +256,7 @@ async fn main() {
                 cpu_batch_size,
                 duration,
                 allow_integrated,
+                gpu_engine,
             )
             .await;
         }
@@ -293,9 +318,9 @@ fn init_logger(verbose: bool) {
     if std::env::var("RUST_LOG").is_err() {
         // Filter out noisy wgpu/naga shader compilation logs
         let log_level = if verbose {
-            "debug,miner=debug,gpu_engine=debug,engine_cpu=debug,wgpu=warn,wgpu_core=warn,wgpu_hal=warn,naga=warn"
+            "debug,miner=debug,gpu_engine=debug,cuda_engine=debug,engine_cpu=debug,wgpu=warn,wgpu_core=warn,wgpu_hal=warn,naga=warn"
         } else {
-            "info,miner=info,gpu_engine=info,wgpu=error,wgpu_core=error,wgpu_hal=error,naga=error"
+            "info,miner=info,gpu_engine=info,cuda_engine=info,wgpu=error,wgpu_core=error,wgpu_hal=error,naga=error"
         };
         std::env::set_var("RUST_LOG", log_level);
     }
@@ -309,6 +334,7 @@ async fn run_benchmark(
     cpu_batch_size: u64,
     duration: u64,
     allow_integrated: bool,
+    gpu_engine_pref: String,
 ) {
     let effective_cpu_workers = cpu_workers.unwrap_or_else(num_cpus::get);
 
@@ -318,6 +344,7 @@ async fn run_benchmark(
         gpu_batch_size,
         0,
         allow_integrated,
+        &gpu_engine_pref,
     ) {
         Ok((engine, count)) => (engine, count),
         Err(e) => {

crates/miner-service/Cargo.toml

@@ -31,4 +31,5 @@ quantus-miner-api = { workspace = true }
 pow-core = { path = "../pow-core" }
 engine-cpu = { path = "../engine-cpu", optional = true }
 engine-gpu = { path = "../engine-gpu" }
+engine-cuda = { path = "../engine-cuda" } # lair: quantus/miner#3
 metrics = { path = "../metrics" }

crates/miner-service/src/lib.rs

@@ -39,6 +39,9 @@ pub struct ServiceConfig {
     pub gpu_throttle_ms: u64,
     /// Allow integrated GPUs even when discrete GPUs are available
     pub allow_integrated: bool,
+    /// lair: GPU engine preference: "auto" (CUDA when available, else wgpu),
+    /// "cuda" or "wgpu" (quantus/miner#3).
+    pub gpu_engine: String,
 }
 
 /// Engine type for tracking metrics per compute type.
@@ -62,6 +65,8 @@ pub struct WorkerResult {
     pub hash_count: u64,
     /// Whether this worker has finished its range.
     pub completed: bool,
+    /// lair: when the candidate was found (seal latency, quantus/miner#9).
+    pub found_at: Option<std::time::Instant>,
 }
 
 /// A successful mining candidate.
@@ -90,6 +95,8 @@ pub struct MiningJob {
     pub ctx: pow_core::JobContext,
     /// Job ID to detect stale results after job transitions
     pub job_id: u64,
+    /// lair: when the job was issued (pickup latency, quantus/miner#9).
+    created_at: std::time::Instant,
 }
 
 /// Persistent worker thread pool that keeps threads alive between jobs.
@@ -208,6 +215,7 @@ impl WorkerPool {
         let job = MiningJob {
             ctx,
             job_id: new_job_id,
+            created_at: std::time::Instant::now(), // lair
         };
 
         // Dispatch job to all workers using bounded channels (capacity 16).
@@ -340,6 +348,9 @@ fn worker_loop(
             log::debug!("[WORKER {type_str}-{thread_id}] Drained {skipped} stale jobs from queue");
         }
 
+        // lair: issued-to-picked-up; for a busy worker this is the cancel latency.
+        metrics::observe_job_pickup(type_str, job.created_at.elapsed());
+
         // Capture the job's ID for later validation
         let job_id = job.job_id;
         log::debug!("[WORKER {type_str}-{thread_id}] Received job {job_id}");
@@ -398,6 +409,7 @@ fn worker_loop(
                 candidate: None, // Discard the stale candidate
                 hash_count,
                 completed: true,
+                found_at: None,
             });
             continue;
         }
@@ -442,6 +454,7 @@ fn worker_loop(
                     candidate: None,
                     hash_count,
                     completed: true,
+                    found_at: None,
                 });
                 break; // Exit the worker loop
             }
@@ -452,6 +465,7 @@ fn worker_loop(
         };
 
         // Send result (non-blocking to avoid deadlock if receiver is full)
+        let found_at = candidate.as_ref().map(|_| std::time::Instant::now()); // lair
         let _ = result_tx.try_send(WorkerResult {
             thread_id,
             engine_type,
@@ -459,12 +473,14 @@ fn worker_loop(
             candidate,
             hash_count,
             completed: true,
+            found_at,
         });
     }
 
     // Clean up GPU resources on thread exit
     if engine_type == EngineType::Gpu {
         engine_gpu::GpuEngine::clear_worker_resources();
+        engine_cuda::CudaEngine::clear_worker_resources(); // lair
     }
 
     log::debug!("{type_str} worker {thread_id} exited");
@@ -476,12 +492,37 @@ pub fn resolve_gpu_configuration(
     batch_size: u32,
     throttle_ms: u64,
     allow_integrated: bool,
+    gpu_engine: &str,
 ) -> anyhow::Result<(Option<Arc<dyn MinerEngine>>, usize)> {
     // Explicit 0 means no GPU
     if requested_devices == Some(0) {
         return Ok((None, 0));
     }
 
+    // lair: native CUDA engine first unless wgpu was asked for (quantus/miner#3).
+    if gpu_engine != "wgpu" {
+        match engine_cuda::CudaEngine::try_new(batch_size, throttle_ms) {
+            Ok(engine) => {
+                let available = engine.device_count();
+                let count = match requested_devices {
+                    Some(n) if n > available => anyhow::bail!(
+                        "Requested {n} GPU devices but only {available} available (CUDA)"
+                    ),
+                    Some(n) => n,
+                    None => {
+                        log::info!("Auto-detected {available} CUDA device(s)");
+                        available
+                    }
+                };
+                return Ok((Some(Arc::new(engine)), count));
+            }
+            Err(e) if gpu_engine == "cuda" => {
+                anyhow::bail!("CUDA engine requested but unavailable: {e}")
+            }
+            Err(e) => log::info!("CUDA engine unavailable ({e}); using wgpu"),
+        }
+    }
+
     // Try to initialize GPU engine
     let engine = engine_gpu::GpuEngine::try_new(batch_size, throttle_ms, allow_integrated);
     let engine = match engine {
@@ -529,6 +570,7 @@ pub async fn run(config: ServiceConfig) -> anyhow::Result<()> {
         config.gpu_batch_size,
         config.gpu_throttle_ms,
         config.allow_integrated,
+        &config.gpu_engine,
     )?;
 
     // Resolve CPU workers
@@ -562,6 +604,14 @@ pub async fn run(config: ServiceConfig) -> anyhow::Result<()> {
         cpu_workers,
         gpu_devices
     );
+    // lair: expose the effective configuration as labels (quantus/miner#9).
+    metrics::set_config_info(
+        gpu_engine.as_ref().map(|e| e.name()).unwrap_or("cpu"),
+        config.gpu_batch_size,
+        gpu_devices,
+        cpu_workers,
+        config.gpu_throttle_ms,
+    );
 
     if let Some(ref engine) = cpu_engine {
         let name = engine.name();

crates/miner-service/src/quic.rs

@@ -38,6 +38,8 @@ pub async fn connect_and_mine(
 
     let mut reconnect_delay = Duration::from_secs(1);
     const MAX_RECONNECT_DELAY: Duration = Duration::from_secs(30);
+    // lair: time without a connection, counted after the first success.
+    let mut disconnected_since: Option<Instant> = None;
 
     loop {
         log::info!("⛏️ Connecting to node at {}...", node_addr);
@@ -45,6 +47,10 @@ pub async fn connect_and_mine(
         match establish_connection(node_addr, auth_token, tls_cert_sha256).await {
             Ok((connection, send, recv)) => {
                 log::info!("⛏️ Connected to node at {}", node_addr);
+                metrics::record_connected(); // lair
+                if let Some(t) = disconnected_since.take() {
+                    metrics::record_disconnected_time(t.elapsed());
+                }
 
                 let mut authenticated = false;
                 if let Err(e) =
@@ -61,6 +67,8 @@ pub async fn connect_and_mine(
                     }
                     log::info!("⛏️ Connection lost: {}", e);
                 }
+                metrics::record_disconnected(); // lair
+                disconnected_since = Some(Instant::now());
                 // Only clear backoff after the node accepted auth (first NewJob).
                 // connect() already treats explicit "auth failed" as permanent;
                 // this covers any other post-Ready close that looked like success.
@@ -76,6 +84,7 @@ pub async fn connect_and_mine(
                     return Err(e);
                 }
                 log::warn!("⛏️ Failed to connect to node: {}", e);
+                metrics::record_connect_failed(); // lair
             }
         }
 
@@ -145,6 +154,8 @@ async fn handle_connection(
     let mut cpu_hashes: u64 = 0;
     let mut gpu_hashes: u64 = 0;
     let mut result_sent_for_current_job = false;
+    // lair: when the last result went out, for node-attributable idle time.
+    let mut result_sent_at: Option<Instant> = None;
 
     log::info!("⛏️ Waiting for mining jobs from node...");
 
@@ -214,7 +225,15 @@ async fn handle_connection(
                         };
 
                         let msg = MinerMessage::JobResult(result);
-                        send_message_checked(&connection, &mut send, &msg).await?;
+                        if let Err(e) = send_message_checked(&connection, &mut send, &msg).await {
+                            metrics::record_result_send_failed(); // lair
+                            return Err(e);
+                        }
+                        // lair: seal latency and the start of the idle window.
+                        metrics::record_result_submitted(
+                            worker_result.found_at.map(|t| t.elapsed()),
+                        );
+                        result_sent_at = Some(Instant::now());
                     }
                 }
             }
@@ -239,6 +258,10 @@ async fn handle_connection(
                 match msg_result {
                     Ok(MinerMessage::NewJob(request)) => {
                         *authenticated = true;
+                        metrics::record_job_received(); // lair
+                        if let Some(t) = result_sent_at.take() {
+                            metrics::record_job_idle(t.elapsed());
+                        }
                         log::info!(
                             "⛏️ Received job: id={}, hash=0x{}",
                             request.job_id,

deploy/README.md

@@ -0,0 +1,42 @@
+# deploy
+
+Everything the CI deploy of the miner ships or needs (`.gitea/workflows/deploy.yaml`).
+Moved from `lair/quantus` on 2026-09-03 so that one repo owns the deployed miner;
+the fleet's monitoring, GPU power limits and nvidia metrics stay there.
+
+| File | Lands at |
+| --- | --- |
+| `quantus-miner.service` | `/etc/systemd/system/quantus-miner.service` |
+| `quantus-miner.sysusers.conf` | `/etc/sysusers.d/quantus-miner.conf` |
+| `quantus-miner-metrics.xml` | `/etc/firewalld/services/quantus-miner-metrics.xml` |
+| `miner.env.tmpl` | rendered to `/etc/quantus-miner/miner.env` |
+| `infra-setup.sh` | operator-run once per host; installs the `gitea_ci` sudoers |
+
+The miner's credentials are generated by the node and copied host-to-host by the
+deploy on every run. Nothing secret lives here.
+
+## Traps, all previously hit
+
+- **Adding a file or privileged command means re-running `infra-setup.sh`.** The
+  deploy preflights the host's `sudo -n -l` against the script's SUDO block and
+  names the missing paths, instead of failing partway with `sudo: a password is required`.
+- **Deploys must be no-ops when nothing changed.** `push()` uses `rsync -ic`
+  (checksum). The artifact is rebuilt every run so mtimes always differ. Restart
+  only on an itemised content change.
+- **`--node-addr` takes an IP, not a name.** It parses as a Rust `SocketAddr`
+  with no DNS. The deploy resolves the node on the miner host and renders the
+  literal; the `10.x` never enters the repo.
+- **Pipes in validate steps.** The shell is `bash -e -o pipefail`; a reader that
+  exits early (`awk ... exit`, `grep -q`, `head`) SIGPIPEs the writer, exit 141.
+  Capture into a variable and parse with a here-string.
+- **ssh argument quoting.** `run()` is `ssh ... "$@"` and the remote shell
+  re-splits; any argument with a space is passed as one pre-quoted string.
+- **`PrivateDevices=false`** in the unit is deliberate: the miner needs `/dev/nvidia*`.
+- **`systemctl is-active` is not evidence of mining.** A miner that found no
+  adapter, or never got a job, is `active`. Validate asserts `miner_hashes_total`
+  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`.
+- **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

@@ -0,0 +1,135 @@
+#!/usr/bin/env bash
+# One-time provisioning of a mining host for the CI deploy in
+# .gitea/workflows/deploy.yaml (quantus/miner#8). The miner role of
+# lair/quantus's script/infra-setup.sh, moved here so the deployer and its
+# sudoers grant list live in the same repo and cannot drift: the deploy
+# preflights the target's `sudo -n -l` against the SUDO block below.
+#
+# Convention: ~/git/architecture/deployment-gitea-actions.md §1–§2. Run from a
+# workstation with admin (sudo) ssh access to the targets — NOT the gitea_ci
+# account. Idempotent; re-running is a no-op. Skips unreachable hosts.
+#
+#   ./deploy/infra-setup.sh --pubkey ~/.ssh/id_gitea_ci.pub
+#   ./deploy/infra-setup.sh --pubkey ~/.ssh/id_gitea_ci.pub --miner-hosts benjy.hanzalova.internal
+#
+# The runner keypair is NOT generated here. It already exists at
+# ~/.ssh/id_gitea_ci and is shared by every project's RSYNC_SSH_KEY secret and
+# every host's gitea_ci authorized_keys — regenerating it would silently break
+# every other deploy on the fleet.
+#
+# Re-run this whenever the deploy gains a new file to ship or a new privileged
+# command; the preflight fails up front naming the missing paths rather than
+# dying partway through an rsync with "sudo: a password is required".
+set -euo pipefail
+
+ADMIN_USER="${ADMIN_USER:-$USER}"
+# Must agree with the deploy matrix in .gitea/workflows/deploy.yaml.
+MINER_HOSTS="${MINER_HOSTS-benjy.hanzalova.internal quadbrat.hanzalova.internal beast.hanzalova.internal}"
+PUBKEY=""
+
+while [ $# -gt 0 ]; do
+  case "$1" in
+    --pubkey)      PUBKEY="$2"; shift 2 ;;
+    --miner-hosts) MINER_HOSTS="$2"; shift 2 ;;
+    --admin)       ADMIN_USER="$2"; shift 2 ;;
+    *) echo "unknown arg: $1" >&2; exit 2 ;;
+  esac
+done
+[ -n "$PUBKEY" ] && [ -r "$PUBKEY" ] || { echo "--pubkey <file> is required and must be readable" >&2; exit 2; }
+
+provision_miner() {
+  local host="$1"
+  echo "== ${host} (miner) =="
+  if ! ssh -o ConnectTimeout=8 -o BatchMode=yes "${ADMIN_USER}@${host}" true; then
+    echo "  ! unreachable as ${ADMIN_USER} — skipping" >&2
+    return 1
+  fi
+
+  # All privileged work in one remote `sudo bash`. The runner pubkey is the only
+  # dynamic value, passed as $1 (single line, no quoting hazard). The sudoers
+  # `\=` are the required escapes — visudo rejects a bare `=` in a command arg.
+  # The `*` in an rsync line matches rsync's --server arg vector; the trailing
+  # literal destination is what actually bounds the rule.
+  ssh "${ADMIN_USER}@${host}" "sudo bash -seu -- '$(cat "$PUBKEY")'" <<'REMOTE'
+    PUBKEY="$1"
+
+    # /bin/bash, NOT nologin. The deploy runs `ssh gitea_ci@host <command>`, and
+    # a nologin shell refuses that with "This account is currently not
+    # available": the key authenticates, then the command cannot run.
+    if ! getent passwd gitea_ci >/dev/null; then
+      useradd --system --create-home --home-dir /var/lib/gitea_ci \
+              --shell /bin/bash gitea_ci
+      echo "  + created gitea_ci"
+    else
+      echo "  = gitea_ci already present"
+    fi
+    cur=$(getent passwd gitea_ci | cut -d: -f7)
+    if [ "$cur" != /bin/bash ]; then
+      usermod -s /bin/bash gitea_ci
+      echo "  ~ gitea_ci shell was ${cur} — set to /bin/bash so ssh commands can run"
+    fi
+
+    install -d -o gitea_ci -g gitea_ci -m 0700 /var/lib/gitea_ci/.ssh
+    ak=/var/lib/gitea_ci/.ssh/authorized_keys
+    touch "$ak"
+    grep -qxF "$PUBKEY" "$ak" || printf '%s\n' "$PUBKEY" >> "$ak"
+    chown gitea_ci:gitea_ci "$ak"
+    chmod 0600 "$ak"
+    usermod -aG systemd-journal gitea_ci
+
+    # The GPU must be present and enumerable before a miner deploy is worth
+    # attempting; failing here beats a green deploy that mines nothing.
+    if ! command -v nvidia-smi >/dev/null; then
+      echo "  ! nvidia-smi not found — this host has no usable NVIDIA driver" >&2
+      exit 1
+    fi
+    nvidia-smi --query-gpu=name --format=csv,noheader | sed 's/^/  = gpu: /'
+
+    tmp=/etc/sudoers.d/.quantus-miner_gitea_ci.tmp
+    cat > "$tmp" <<'SUDO'
+gitea_ci ALL=(root) NOPASSWD: /usr/bin/rsync * /usr/local/bin/quantus-miner
+gitea_ci ALL=(root) NOPASSWD: /usr/bin/rsync * /etc/systemd/system/quantus-miner.service
+gitea_ci ALL=(root) NOPASSWD: /usr/bin/rsync * /etc/sysusers.d/quantus-miner.conf
+gitea_ci ALL=(root) NOPASSWD: /usr/bin/rsync * /etc/quantus-miner/miner.env
+gitea_ci ALL=(root) NOPASSWD: /usr/bin/rsync * /etc/quantus-miner/miner-auth-token
+gitea_ci ALL=(root) NOPASSWD: /usr/bin/rsync * /etc/quantus-miner/miner-tls-cert-sha256
+gitea_ci ALL=(root) NOPASSWD: /usr/bin/rsync * /etc/firewalld/services/quantus-miner-metrics.xml
+gitea_ci ALL=(root) NOPASSWD: /usr/bin/firewall-cmd --get-default-zone
+gitea_ci ALL=(root) NOPASSWD: /usr/bin/firewall-cmd --reload
+gitea_ci ALL=(root) NOPASSWD: /usr/bin/firewall-cmd --permanent --zone\=* --add-rich-rule\=*
+gitea_ci ALL=(root) NOPASSWD: /usr/bin/firewall-cmd --zone\=* --add-rich-rule\=*
+gitea_ci ALL=(root) NOPASSWD: /usr/bin/firewall-cmd --zone\=* --query-rich-rule\=*
+gitea_ci ALL=(root) NOPASSWD: /usr/bin/systemd-sysusers
+gitea_ci ALL=(root) NOPASSWD: /usr/bin/install -d -o root -g quantus-miner -m 0750 /etc/quantus-miner
+gitea_ci ALL=(root) NOPASSWD: /usr/bin/install -d -o quantus-miner -g quantus-miner -m 0750 /var/lib/quantus-miner
+gitea_ci ALL=(root) NOPASSWD: /usr/sbin/restorecon -R /usr/local/bin/quantus-miner /etc/quantus-miner /var/lib/quantus-miner
+gitea_ci ALL=(root) NOPASSWD: /usr/bin/systemctl daemon-reload
+gitea_ci ALL=(root) NOPASSWD: /usr/bin/systemctl enable quantus-miner.service
+gitea_ci ALL=(root) NOPASSWD: /usr/bin/systemctl restart quantus-miner.service
+# stop/start (not just restart) so the benchmark harness (bench.yaml) can pause
+# 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
+# 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;
+# this is what makes that survivable.
+gitea_ci ALL=(root) NOPASSWD: /usr/bin/cp -p /usr/local/bin/quantus-miner /usr/local/bin/quantus-miner.prev
+gitea_ci ALL=(root) NOPASSWD: /usr/bin/install -m 0755 /usr/local/bin/quantus-miner.prev /usr/local/bin/quantus-miner
+SUDO
+    chmod 0440 "$tmp"
+    visudo -cf "$tmp"
+    mv "$tmp" /etc/sudoers.d/quantus-miner_gitea_ci
+    echo "  = sudoers quantus-miner_gitea_ci installed and visudo-verified"
+REMOTE
+}
+
+rc=0
+for h in $MINER_HOSTS; do
+  provision_miner "$h" || rc=1
+done
+echo
+echo "done. remaining operator steps:"
+echo "  1. confirm RSYNC_SSH_KEY is set in the repo's Actions secrets"
+echo "  2. push to main, or run the deploy workflow manually (mode: validate first)"
+exit $rc

deploy/miner.env.tmpl

@@ -0,0 +1,6 @@
+# Rendered by .gitea/workflows/deploy.yaml (quantus/miner) and rsynced to
+# /etc/quantus-miner/miner.env (0640 root:quantus-miner). Not secret — the
+# miner's actual credentials are the auth token and TLS pin, which are copied
+# from the node host as separate 0640 files.
+QUANTUS_NODE_ADDR={{QUANTUS_NODE_ADDR}}
+QUANTUS_GPU_DEVICES={{QUANTUS_GPU_DEVICES}}

deploy/quantus-miner-metrics.xml

@@ -0,0 +1,9 @@
+<?xml version="1.0" encoding="utf-8"?>
+<service>
+  <short>quantus-miner-metrics</short>
+  <description>Quantus miner Prometheus exporter. The miner binds this on
+  0.0.0.0 unconditionally (crates/metrics: SocketAddr from [0,0,0,0]) — there is
+  no loopback option — so the firewalld rich rule scoped to the scrape host is
+  the only thing bounding who can reach it.</description>
+  <port protocol="tcp" port="9900"/>
+</service>

deploy/quantus-miner.service

@@ -0,0 +1,64 @@
+# Quantus external miner. Runs on the GPU host and connects OUT to the node's
+# QUIC control channel; it listens on nothing but its loopback metrics port, so
+# it ships no firewalld service of its own.
+#
+# Hardened per ~/git/architecture/generic.md §8, with one deliberate relaxation:
+#
+#   PrivateDevices=false
+#     The miner needs /dev/nvidia*; PrivateDevices=true hides them and wgpu
+#     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.
+#
+# --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.
+
+[Unit]
+Description=Quantus external miner (GPU)
+Documentation=https://github.com/Quantus-Network/quantus-miner
+After=network-online.target
+Wants=network-online.target
+
+[Service]
+Type=simple
+User=quantus-miner
+Group=quantus-miner
+Environment=RUST_LOG=info
+WorkingDirectory=/var/lib/quantus-miner
+
+ExecStart=/usr/local/bin/quantus-miner serve \
+    --node-addr ${QUANTUS_NODE_ADDR} \
+    --auth-token-file /etc/quantus-miner/miner-auth-token \
+    --tls-cert-sha256-file /etc/quantus-miner/miner-tls-cert-sha256 \
+    --gpu-devices ${QUANTUS_GPU_DEVICES} \
+    --cpu-workers 0 \
+    --metrics-port 9900
+EnvironmentFile=/etc/quantus-miner/miner.env
+
+# The node may be down, still syncing, or mid-redeploy; reconnecting is normal
+# operation, not an error condition.
+Restart=always
+RestartSec=15s
+LimitNOFILE=65536
+
+NoNewPrivileges=true
+ProtectSystem=strict
+ProtectHome=true
+PrivateTmp=true
+PrivateDevices=false
+ProtectKernelTunables=true
+ProtectKernelModules=true
+ProtectControlGroups=true
+RestrictRealtime=true
+RestrictSUIDSGID=true
+LockPersonality=true
+SystemCallArchitectures=native
+MemoryDenyWriteExecute=false
+
+ReadWritePaths=/var/lib/quantus-miner
+RestrictAddressFamilies=AF_UNIX AF_INET AF_INET6
+
+[Install]
+WantedBy=multi-user.target

deploy/quantus-miner.sysusers.conf

@@ -0,0 +1,4 @@
+#Type Name            ID  GECOS                           Home directory         Shell
+u     quantus-miner   -   "Quantus miner service account" /var/lib/quantus-miner /usr/sbin/nologin
+m     quantus-miner   video
+m     quantus-miner   render