Merge pull request 'deploy: require nvcc at build time' (#32) from ci/deploy-require-cuda into main

The CUDA build script warns and embeds an empty kernel when nvcc is
absent, which is right for the rust runner's test job and wrong for a
deploy build. MINER_CUDA_REQUIRE=1 turns that into a build failure
before any host is touched; the release workflow already sets it.

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

.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,397 @@
+---
+# 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
+      - .gitea/workflows/release.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 }}
+          # Fail the build here if nvcc is missing rather than shipping a
+          # binary that falls back to wgpu and failing validate on the hosts.
+          # (The same warning in ci.yml's test job is expected: the `rust`
+          # runner has no nvcc and only lints and tests.)
+          MINER_CUDA_REQUIRE: "1"
+        run: |
+          set -euo pipefail
+          cargo build --release --locked -p miner-cli
+          ./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.
+          # `cuda_kernel` selects the kernel entry (MINER_CUDA_KERNEL, #29):
+          # unrolled is the default; loop measured +1.7% on the 4090 and
+          # -1.8% / -3.4% on the 3060 / 5090 (2026-09-14), so only benjy runs
+          # it. `kernel` must be the id that choice reports: cuda or cuda-loop.
+          - host: benjy.hanzalova.internal
+            node: bob.hanzalova.internal
+            gpu_devices: "1"        # 1x RTX 4090 (sm_89); reference card for #2
+            cuda_kernel: loop
+            kernel: cuda-loop
+          - host: quadbrat.hanzalova.internal
+            node: bob.hanzalova.internal
+            gpu_devices: "1"        # 1x RTX 3060 (sm_86)
+            cuda_kernel: unrolled
+            kernel: cuda
+          - host: beast.hanzalova.internal
+            node: bob.hanzalova.internal
+            gpu_devices: "2"        # 2x RTX 5090 (sm_120)
+            cuda_kernel: unrolled
+            kernel: cuda
+    steps:
+      - uses: actions/checkout@v4
+      - 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 }}
+          CUDA_KERNEL: ${{ matrix.cuda_kernel }}
+        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"])
+          t = t.replace("{{MINER_CUDA_KERNEL}}", os.environ["CUDA_KERNEL"])
+          pathlib.Path("miner.env").write_text(t)
+          PY
+          push --chown=root:quantus-miner --chmod=F0640 \
+            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

.gitea/workflows/release.yaml

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

.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,91 +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: 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"
@@ -577,7 +603,7 @@ dependencies = [
 
 [[package]]
 name = "engine-cpu"
-version = "4.0.1"
+version = "4.0.2"
 dependencies = [
  "criterion",
  "hex",
@@ -586,9 +612,22 @@ 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.1"
+version = "4.0.2"
 dependencies = [
  "bytemuck",
  "criterion",
@@ -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"
@@ -1476,7 +1526,7 @@ dependencies = [
 
 [[package]]
 name = "metrics"
-version = "4.0.1"
+version = "4.0.2"
 dependencies = [
  "anyhow",
  "log",
@@ -1504,7 +1554,7 @@ dependencies = [
 
 [[package]]
 name = "miner-cli"
-version = "4.0.1"
+version = "4.0.2"
 dependencies = [
  "clap",
  "engine-cpu",
@@ -1522,11 +1572,12 @@ dependencies = [
 
 [[package]]
 name = "miner-service"
-version = "4.0.1"
+version = "4.0.2"
 dependencies = [
  "anyhow",
  "crossbeam-channel",
  "engine-cpu",
+ "engine-cuda",
  "engine-gpu",
  "getrandom 0.2.17",
  "hex",
@@ -1543,7 +1594,7 @@ dependencies = [
 
 [[package]]
 name = "miner-telemetry"
-version = "4.0.1"
+version = "4.0.2"
 dependencies = [
  "anyhow",
  "futures",
@@ -2003,7 +2054,7 @@ dependencies = [
 
 [[package]]
 name = "pool-service"
-version = "4.0.1"
+version = "4.0.2"
 dependencies = [
  "anyhow",
  "clap",
@@ -2049,7 +2100,7 @@ dependencies = [
 
 [[package]]
 name = "pow-core"
-version = "4.0.1"
+version = "4.0.2"
 dependencies = [
  "hex",
  "primitive-types 0.13.1",
@@ -2291,7 +2342,7 @@ dependencies = [
 
 [[package]]
 name = "quic-transport"
-version = "4.0.1"
+version = "4.0.2"
 dependencies = [
  "anyhow",
  "quantus-miner-api",
@@ -2871,7 +2922,7 @@ dependencies = [
 
 [[package]]
 name = "solver-wasm"
-version = "4.0.1"
+version = "4.0.2"
 dependencies = [
  "primitive-types 0.13.1",
  "qpow-math",
@@ -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",
@@ -18,7 +20,7 @@ resolver = "2"
 edition = "2021"
 authors = ["Quantus Network"]
 description = "Quantus External Miner Workspace"
-version = "4.0.1"
+version = "4.0.2"
 
 [workspace.dependencies]
 anyhow = "1"

README.md

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

crates/bench-harness/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,18 @@
+[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"] }
+qp-poseidon-constants = "1.1"
+
+[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/lib.rs

@@ -0,0 +1,681 @@
+#![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");
+/// lair: which kernel entry the module runs (quantus/miner#27). `unrolled`
+/// is `mining_main`, the deployed kernel; `loop` is `mining_loop`, the
+/// rolled-round twin under evaluation. Selected with MINER_CUDA_KERNEL so a
+/// host can be switched without a rebuild, and reported as the kernel id on
+/// the per-device metrics so the deploy validate and the dashboards can tell
+/// them apart.
+#[derive(Clone, Copy, PartialEq, Eq, Debug)]
+enum KernelChoice {
+    Unrolled,
+    Loop,
+}
+
+impl KernelChoice {
+    fn from_env() -> Self {
+        match std::env::var("MINER_CUDA_KERNEL").as_deref() {
+            Ok("loop") => KernelChoice::Loop,
+            Ok("unrolled") | Ok("") | Err(_) => KernelChoice::Unrolled,
+            Ok(other) => {
+                log::warn!(target: "cuda_engine", "MINER_CUDA_KERNEL={other:?} is not unrolled|loop; using unrolled");
+                KernelChoice::Unrolled
+            }
+        }
+    }
+    fn function(self) -> &'static str {
+        match self {
+            KernelChoice::Unrolled => "mining_main",
+            KernelChoice::Loop => "mining_loop",
+        }
+    }
+    /// Kernel id on metrics: `cuda` for the deployed kernel (what the deploy
+    /// validate asserts), `cuda-loop` for the candidate.
+    fn id(self) -> &'static str {
+        match self {
+            KernelChoice::Unrolled => "cuda",
+            KernelChoice::Loop => "cuda-loop",
+        }
+    }
+}
+/// Threads per block. Must match the kernel's `__launch_bounds__` (LAIR_TPB,
+/// default 256); override for experiments with MINER_CUDA_THREADS_PER_BLOCK
+/// on a kernel compiled with the same value.
+const THREADS_PER_BLOCK_DEFAULT: u32 = 256;
+/// The loop kernel's `LAIR_LOOP_TPB`.
+const LOOP_THREADS_PER_BLOCK_DEFAULT: u32 = 512;
+
+fn threads_per_block(kernel: KernelChoice) -> u32 {
+    std::env::var("MINER_CUDA_THREADS_PER_BLOCK")
+        .ok()
+        .and_then(|v| v.parse().ok())
+        .filter(|v: &u32| *v > 0 && v.is_multiple_of(32))
+        .unwrap_or(match kernel {
+            KernelChoice::Unrolled => THREADS_PER_BLOCK_DEFAULT,
+            KernelChoice::Loop => LOOP_THREADS_PER_BLOCK_DEFAULT,
+        })
+}
+
+/// lair: size every batch to whole grids. The CLI default batch (1M nonces)
+/// fills a 5090 about eleven waves deep with a half-empty tail wave and pays
+/// the per-launch gap every 1M hashes. Rounding the batch up to a multiple of
+/// the grid (SMs x threads per SM, one nonce per thread) measured +1.6% on
+/// beast, and two grids per batch (each thread loops over two nonces) a
+/// further +0.4% there and +0.5% on the 4090, neutral on the 3060
+/// (quantus/miner#3, 2026-09-13). A batch is then ~10 ms on every card, so a
+/// job switch discards about 0.1% of a block interval. Overrides:
+/// MINER_CUDA_BATCH_ALIGN=0 for the CLI batch as given, MINER_CUDA_BATCH_WAVES
+/// for the minimum number of grids per batch.
+const BATCH_WAVES_DEFAULT: u32 = 2;
+
+fn batch_align() -> bool {
+    std::env::var("MINER_CUDA_BATCH_ALIGN")
+        .map(|v| v != "0")
+        .unwrap_or(true)
+}
+
+fn batch_waves() -> u32 {
+    std::env::var("MINER_CUDA_BATCH_WAVES")
+        .ok()
+        .and_then(|v| v.parse().ok())
+        .filter(|v: &u32| *v > 0)
+        .unwrap_or(BATCH_WAVES_DEFAULT)
+}
+/// 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,
+    kernel: KernelChoice,
+    threads_per_block: 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],
+    /// lair: the state after the first external linear layer (and its round
+    /// constant) for the batch's first nonce. That layer is linear, and only
+    /// the least significant nonce limb changes within a batch, so the kernel
+    /// derives each nonce's state from this by one scalar times a fixed
+    /// small-integer column instead of recomputing the layer (quantus/miner#3).
+    layer0_base: [u64; 12],
+    /// lair: the internal layer's diagonal, for the loop kernel (#27). In the
+    /// parameter bank ptxas reads it straight from the instruction operand;
+    /// from `__constant__` it reloaded all twelve values every round.
+    diag: [u64; 12],
+}
+
+// SAFETY: `#[repr(C)]` plain-old-data with no padding (12 x u64, 32 x u32, 24 x u64),
+// matching the kernel's parameter layout; there is nothing to validate beyond
+// the layout, which is what this marker asserts.
+unsafe impl DeviceRepr for MiningUniforms {}
+
+// 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>() == 416);
+
+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 [sm_{{{}}}]",
+                        d.name,
+                        d.sm_count,
+                        d.kernel.id(),
+                        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 kernel = KernelChoice::from_env();
+        let func = module.load_function(kernel.function())?;
+        Ok(Self {
+            ctx,
+            func,
+            name,
+            sm_count,
+            kernel,
+            threads_per_sm: threads_per_sm(),
+            threads_per_block: threads_per_block(kernel),
+            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 want: u32 = if batch_align() {
+                let grid = (dev.sm_count * dev.threads_per_sm).max(1);
+                let waves = self.batch_size.div_ceil(grid).max(batch_waves());
+                waves.saturating_mul(grid)
+            } else {
+                self.batch_size
+            };
+            let this_batch: u32 = if cap > U512::from(want) {
+                want
+            } 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 tpb = dev.threads_per_block;
+        let num_blocks = ((logical_threads as u32).div_ceil(tpb)).max(1);
+        let total_threads = (num_blocks * tpb) 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 layer0_base = first_layer_after_absorb(&midstate, &start_u32s[..8]);
+        let uni = MiningUniforms {
+            midstate,
+            start_nonce: start_u32s,
+            target: res.target_u32s,
+            layer0_base,
+            diag: qp_poseidon_constants::POSEIDON2_MATRIX_DIAG_12_RAW,
+        };
+
+        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: (tpb, 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()
+}
+
+// lair: host-side twin of the kernel's absorb + first external linear layer,
+// in canonical Goldilocks arithmetic. Bit-exact with `ext_layer_rc` in
+// kernels/mining.cu modulo p (the kernel keeps lazy representatives).
+const GOLDILOCKS_P: u128 = 0xFFFF_FFFF_0000_0001;
+
+fn gf_add(a: u64, b: u64) -> u64 {
+    ((a as u128 + b as u128) % GOLDILOCKS_P) as u64
+}
+
+/// State after absorbing the low nonce limbs into the midstate and applying
+/// the first external layer plus the first round constant, for the batch's
+/// first nonce. `nonce_le` are the low 8 little-endian u32 limbs.
+fn first_layer_after_absorb(midstate: &[u64; 12], nonce_le: &[u32]) -> [u64; 12] {
+    let mut st = *midstate;
+    for i in 0..8 {
+        st[i] = gf_add(st[i], nonce_le[7 - i].swap_bytes() as u64);
+    }
+    let mut out = [0u64; 12];
+    for chunk in 0..3 {
+        let o = chunk * 4;
+        let (x0, x1, x2, x3) = (st[o], st[o + 1], st[o + 2], st[o + 3]);
+        let t01 = gf_add(x0, x1);
+        let t23 = gf_add(x2, x3);
+        let t0123 = gf_add(t01, t23);
+        let t01123 = gf_add(t0123, x1);
+        let t01233 = gf_add(t0123, x3);
+        out[o + 3] = gf_add(t01233, gf_add(x0, x0));
+        out[o + 1] = gf_add(t01123, gf_add(x2, x2));
+        out[o] = gf_add(t01123, t01);
+        out[o + 2] = gf_add(t01233, t23);
+    }
+    let mut sums = [0u64; 4];
+    for k in 0..4 {
+        sums[k] = gf_add(gf_add(out[k], out[k + 4]), out[k + 8]);
+    }
+    let rc = &qp_poseidon_constants::POSEIDON2_INITIAL_EXTERNAL_CONSTANTS_RAW[0];
+    for i in 0..12 {
+        out[i] = gf_add(gf_add(out[i], sums[i & 3]), rc[i]);
+    }
+    out
+}
+
+#[cfg(test)]
+mod layer0_tests {
+    use super::*;
+
+    /// The kernel's fast path assumes column 7 of the external matrix is this
+    /// vector: applying the layer to the unit vector e7 must reproduce it.
+    #[test]
+    fn column_seven_of_external_matrix() {
+        let zero = [0u64; 12];
+        // absorb puts bswap(limb 0) into element 7: choose limb 0 so that it becomes 1
+        let mut nonce = [0u32; 8];
+        nonce[0] = 1u32.swap_bytes();
+        let with = first_layer_after_absorb(&zero, &nonce);
+        let without = first_layer_after_absorb(&zero, &[0u32; 8]);
+        let rc = &qp_poseidon_constants::POSEIDON2_INITIAL_EXTERNAL_CONSTANTS_RAW[0];
+        let expect: [u64; 12] = [1, 1, 3, 2, 2, 2, 6, 4, 1, 1, 3, 2];
+        for i in 0..12 {
+            assert_eq!(without[i], rc[i]);
+            let diff = (with[i] as u128 + GOLDILOCKS_P - without[i] as u128) % GOLDILOCKS_P;
+            assert_eq!(diff as u64, expect[i], "column entry {i}");
+        }
+    }
+}

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/examples/trusted_hashrate.rs

@@ -46,9 +46,10 @@ impl Runner {
             .await
             .unwrap();
 
+        let kernel = engine_gpu::Kernel::for_adapter(&adapter);
         let desc = wgpu::ShaderModuleDescriptor {
-            label: None,
-            source: wgpu::ShaderSource::Wgsl(include_str!("../src/mining_u64.wgsl").into()),
+            label: Some(kernel.label()),
+            source: wgpu::ShaderSource::Wgsl(kernel.source().into()),
         };
         let shader = if trusted {
             unsafe {

crates/engine-gpu/src/kernels/mining_u64.wgsl

@@ -0,0 +1,573 @@
+// Native-u64 Poseidon2 mining kernel. Default path for non-Apple GPUs.
+// Requires wgpu Features::SHADER_INT64. Bit-exact with mining.wgsl / pow_core.
+
+@group(0) @binding(0) var<storage, read_write> results: array<atomic<u32>>;
+// Sponge state after absorbing header + high nonce half (12 felts as LE u32 pairs),
+// precomputed on the host per batch. See pow_core::mining_midstate.
+@group(0) @binding(1) var<storage, read> midstate: array<u32, 24>;
+@group(0) @binding(2) var<storage, read> start_nonce: array<u32, 16>;
+@group(0) @binding(3) var<storage, read> difficulty_target: array<u32, 16>;
+@group(0) @binding(4) var<storage, read> dispatch_config: array<u32, 3>;
+
+const P64: u64 = 0xFFFFFFFF00000001lu;
+// EPS64 = 2^32 - 1 = 2^64 mod P
+const EPS64: u64 = 0xFFFFFFFFlu;
+
+const RC_INTERNAL: array<u64, 22> = array<u64, 22>(
+    0x97f7798a784ad863lu, 0xd1d2bf082f60d4f0lu, 0x69a377a79f9ad206lu, 0xa9d06906a3858e24lu, 0x295275001eede5b5lu, 0x5874e441117bd746lu, 0x8a084bbba8ed86cclu, 0x3defd7645cde6425lu, 0x3998cfe6871cc137lu, 0x3e52ef8bca48314alu, 0x964a209f85dc9ecclu, 0x3fcc9ee82cc4577elu, 0x8e79b4a5d0096d6dlu, 0x8492362ad2392556lu, 0xee72f470262574d6lu, 0x1e0e18496da2444alu, 0x0f3a74bf215eaac6lu, 0x1b061b76a1c0ded3lu, 0x192c42d86803d7a6lu, 0xf6d49ff997ae0260lu, 0x3ec372e7a0fa3786lu, 0x5538cdf4f23445d3lu
+);
+const RC_INITIAL: array<array<u64, 12>, 4> = array<array<u64, 12>, 4>(
+    array<u64, 12>(0xc002e770975b1607lu, 0xbca51a8dfe14593alu, 0x72938dfbe774f7f9lu, 0xe4f2fe29e03234aclu, 0xd5e0ba2f541b6449lu, 0xec33b868f3cc46c1lu, 0x486dcb55419d475alu, 0x6c1cb2a358cc24f1lu, 0xe3f30d509a1436bblu, 0xd9a64f068dca7c29lu, 0xe59b3f57aabba1aelu, 0x2a3dd4505b478fdclu),
+    array<u64, 12>(0xada1f8dc7676ed25lu, 0x2711aa8b5509d516lu, 0x4ae6acd0c9c92897lu, 0x56eb3d6b5256d67alu, 0x1f7a9d55923bf51elu, 0x3600427d397a7f68lu, 0xe5076df75b72c3d0lu, 0xfcd59aa12c6090adlu, 0xcd895e8c68b57a9elu, 0x41df7ef9d730ae3elu, 0xee3e2b889abe977dlu, 0xd29bb7edbeb9c405lu),
+    array<u64, 12>(0x7d5c08eef608e382lu, 0x89ae889caaf0802clu, 0xb35a8e976d2af617lu, 0xdb14234eafaf5173lu, 0x78f04462d48b1c98lu, 0x265293b0e47ce88alu, 0x999a649b69b9d32flu, 0x64b0a186698e01d3lu, 0xee0b22d0dfae8bb8lu, 0x4fd53e50ca04a7eelu, 0x5762bfe181f25047lu, 0xf51593e2beb5e3bdlu),
+    array<u64, 12>(0x1e5e2b5760e32477lu, 0x622462a1f9aaaeedlu, 0xaa284b3ecdb222aelu, 0x63c8e72f542bf3fclu, 0x3ba588cacb43b5e0lu, 0x23eda6f3c99150ddlu, 0xaad3bea4baac9a5alu, 0xe9da8d699b94184alu, 0xcdb13f4cd93e024clu, 0x902cbd0956f655e3lu, 0x5b4e40ffc759532flu, 0xde795c20a2357af7lu)
+);
+const RC_TERMINAL: array<array<u64, 12>, 4> = array<array<u64, 12>, 4>(
+    array<u64, 12>(0x7b72c539e0ea4c6elu, 0x144573dae2ce9976lu, 0x802028b68f35fc88lu, 0x6d36c5022c4fe7c2lu, 0xa205d0ffa9b9def3lu, 0xf6e7e38b1ea6ba2flu, 0x34f7909ae5258d64lu, 0xb0464d9d77b97fcalu, 0x64ddb9d5de7e00a6lu, 0x0ed0d75c27975d97lu, 0x1cbb36f11127338blu, 0x6673e505cfd0b6balu),
+    array<u64, 12>(0x605f902830872e01lu, 0x3fd5eb927e95fe4flu, 0xe81025b5a24c69cdlu, 0xf7d0ce75de23f74elu, 0xf39942b6a8585089lu, 0x6d808a08f7b71df6lu, 0xf8806b6588f49a8blu, 0x57df2d8c2a32107alu, 0x16e7c2074d654a2dlu, 0x213de241fcf33835lu, 0xb0f2b8905a0976f6lu, 0xd8e3cf2bbd355417lu),
+    array<u64, 12>(0xe498691679d9330flu, 0x763b45d2a3821b28lu, 0x0908bf65eb0a1f0dlu, 0x7691eb2d194b24f4lu, 0x0e43551233ae13b2lu, 0x93c393dbfc2fe76flu, 0x98f607485d48cdealu, 0xe3d95f30309819c0lu, 0x1ef581a93eaf6acflu, 0x0b24c1b7a030fca4lu, 0x624370be5670b327lu, 0x5f1e28615a11e486lu),
+    array<u64, 12>(0xfe04051f909e042blu, 0x7257e5b147fd3803lu, 0xe6ae134bb82f2e78lu, 0x5711fd5cf4784511lu, 0xf83a42660c08c0bclu, 0x2cd8c96d9a3ce855lu, 0x7d2ffb1bb0e17271lu, 0x85ae1528caea3811lu, 0x52a345d5c7adb0b8lu, 0x504c4c51f3faee94lu, 0xbce34a649cfccaf9lu, 0xe0a3389266fb6dc9lu)
+);
+const MDS_DIAG: array<u64, 12> = array<u64, 12>(
+    0xc3b6c08e23ba9300lu, 0xd84b5de94a324fb6lu, 0x0d0c371c5b35b84flu, 0x7964f570e7188037lu, 0x5daf18bbd996604blu, 0x6743bc47b9595257lu, 0x5528b9362c59bb70lu, 0xac45e25b7127b68blu, 0xa2077d7dfbb606b5lu, 0xf3faac6faee378aelu, 0x0c6388b51545e883lu, 0xd27dbb6944917b60lu
+);
+
+// a + b mod P in lazy form. Wrapping carries fold back via 2^64 ≡ EPS64 (mod P).
+fn gf64_add(a: u64, b: u64) -> u64 {
+    let s0 = a + b;
+    let c1 = s0 < a;
+    let s1 = s0 + select(0lu, EPS64, c1);
+    let c2 = c1 && (s1 < s0);
+    return s1 + select(0lu, EPS64, c2);
+}
+
+// Reduce a 128-bit value (lo + hi*2^64) mod P using
+// 2^64 ≡ EPS64 and 2^96 ≡ -1 (mod P).
+fn gf64_reduce(lo: u64, hi: u64) -> u64 {
+    let hi_hi = hi >> 32u;
+    let hi_lo = hi & EPS64;
+    var t0 = lo - hi_hi;
+    t0 = t0 - select(0lu, EPS64, lo < hi_hi);
+    let t1 = hi_lo * EPS64;
+    let t2 = t0 + t1;
+    return t2 + select(0lu, EPS64, t2 < t0);
+}
+
+fn gf64_mul(a: u64, b: u64) -> u64 {
+    let a_lo = a & EPS64;
+    let a_hi = a >> 32u;
+    let b_lo = b & EPS64;
+    let b_hi = b >> 32u;
+    let ll = a_lo * b_lo;
+    let lh = a_lo * b_hi;
+    let hl = a_hi * b_lo;
+    let hh = a_hi * b_hi;
+    let mid = lh + hl;
+    let mid_c = select(0lu, 1lu, mid < lh);
+    let lo = ll + (mid << 32u);
+    let lo_c = select(0lu, 1lu, lo < ll);
+    let hi = hh + (mid >> 32u) + (mid_c << 32u) + lo_c;
+    return gf64_reduce(lo, hi);
+}
+
+fn gf64_sqr(a: u64) -> u64 {
+    let a_lo = a & EPS64;
+    let a_hi = a >> 32u;
+    let ll = a_lo * a_lo;
+    let lh = a_lo * a_hi;
+    let hh = a_hi * a_hi;
+    let mid = lh << 1u;
+    let mid_c = lh >> 63u;
+    let lo = ll + (mid << 32u);
+    let lo_c = select(0lu, 1lu, lo < ll);
+    let hi = hh + (mid >> 32u) + (mid_c << 32u) + lo_c;
+    return gf64_reduce(lo, hi);
+}
+
+fn gf64_sbox(x: u64) -> u64 {
+    let x2 = gf64_sqr(x);
+    let x4 = gf64_sqr(x2);
+    let x6 = gf64_mul(x4, x2);
+    return gf64_mul(x6, x);
+}
+
+fn gf64_canon(a: u64) -> u64 {
+    return a - select(0lu, P64, a >= P64);
+}
+
+// External linear layer: 4x4 MDS on each chunk, then circulant sums.
+fn ext_layer64(state: ptr<function, array<u64, 12>>) {
+    for (var chunk = 0u; chunk < 3u; chunk++) {
+        let o = chunk * 4u;
+        let x0 = (*state)[o];
+        let x1 = (*state)[o + 1u];
+        let x2 = (*state)[o + 2u];
+        let x3 = (*state)[o + 3u];
+        let t01 = gf64_add(x0, x1);
+        let t23 = gf64_add(x2, x3);
+        let t0123 = gf64_add(t01, t23);
+        let t01123 = gf64_add(t0123, x1);
+        let t01233 = gf64_add(t0123, x3);
+        (*state)[o + 3u] = gf64_add(t01233, gf64_add(x0, x0));
+        (*state)[o + 1u] = gf64_add(t01123, gf64_add(x2, x2));
+        (*state)[o] = gf64_add(t01123, t01);
+        (*state)[o + 2u] = gf64_add(t01233, t23);
+    }
+    var sums: array<u64, 4>;
+    for (var k = 0u; k < 4u; k++) {
+        sums[k] = gf64_add(gf64_add((*state)[k], (*state)[k + 4u]), (*state)[k + 8u]);
+    }
+    for (var i = 0u; i < 12u; i++) {
+        (*state)[i] = gf64_add((*state)[i], sums[i % 4u]);
+    }
+}
+
+// Internal linear layer: diagonal matrix plus full sum.
+fn int_layer64(state: ptr<function, array<u64, 12>>) {
+    var sum = (*state)[0];
+    for (var i = 1u; i < 12u; i++) {
+        sum = gf64_add(sum, (*state)[i]);
+    }
+    for (var i = 0u; i < 12u; i++) {
+        (*state)[i] = gf64_add(gf64_mul((*state)[i], MDS_DIAG[i]), sum);
+    }
+}
+
+fn permute64(state: ptr<function, array<u64, 12>>) {
+    ext_layer64(state);
+    for (var r = 0u; r < 4u; r++) {
+        for (var i = 0u; i < 12u; i++) {
+            (*state)[i] = gf64_add((*state)[i], RC_INITIAL[r][i]);
+        }
+        for (var i = 0u; i < 12u; i++) {
+            (*state)[i] = gf64_sbox((*state)[i]);
+        }
+        ext_layer64(state);
+    }
+    for (var r = 0u; r < 22u; r++) {
+        (*state)[0] = gf64_sbox(gf64_add((*state)[0], RC_INTERNAL[r]));
+        int_layer64(state);
+    }
+    for (var r = 0u; r < 4u; r++) {
+        for (var i = 0u; i < 12u; i++) {
+            (*state)[i] = gf64_add((*state)[i], RC_TERMINAL[r][i]);
+        }
+        for (var i = 0u; i < 12u; i++) {
+            (*state)[i] = gf64_sbox((*state)[i]);
+        }
+        ext_layer64(state);
+    }
+}
+
+fn bswap32(v: u32) -> u32 {
+    return ((v & 0xFFu) << 24u) | ((v & 0xFF00u) << 8u) | ((v >> 8u) & 0xFF00u) | (v >> 24u);
+}
+
+@compute @workgroup_size(256)
+fn mining_main(@builtin(global_invocation_id) global_id: vec3<u32>) {
+    if (atomicLoad(&results[0]) != 0u) {
+        return;
+    }
+    let thread_id = global_id.x;
+    let total_threads = dispatch_config[0];
+    let nonces_per_thread = dispatch_config[1];
+    let total_nonces = dispatch_config[2];
+    if (thread_id >= total_threads) {
+        return;
+    }
+    let base_index = thread_id * nonces_per_thread;
+
+    // Hoist uniform storage reads out of the nonce loop
+    var mid: array<u64, 12>;
+    for (var i = 0u; i < 12u; i++) {
+        mid[i] = (u64(midstate[2u * i + 1u]) << 32u) | u64(midstate[2u * i]);
+    }
+    var tgt: array<u32, 16>;
+    for (var i = 0u; i < 16u; i++) {
+        tgt[i] = difficulty_target[i];
+    }
+    var nonce_base: array<u32, 16>;
+    for (var i = 0u; i < 16u; i++) {
+        nonce_base[i] = start_nonce[i];
+    }
+
+    for (var j = 0u; j < nonces_per_thread; j = j + 1u) {
+        let logical_index = base_index + j;
+        if (logical_index >= total_nonces) {
+            break;
+        }
+        if (j > 0u && atomicLoad(&results[0]) != 0u) {
+            return;
+        }
+
+        // The host guarantees a batch never carries into the high nonce half
+        // (limbs 8..15), so only the low 256 bits are incremented here.
+        var current_nonce: array<u32, 16>;
+        let val0 = nonce_base[0];
+        let sum0 = val0 + logical_index;
+        current_nonce[0] = sum0;
+        var carry = select(0u, 1u, sum0 < val0);
+        for (var i = 1u; i < 8u; i++) {
+            let val = nonce_base[i];
+            let sum = val + carry;
+            current_nonce[i] = sum;
+            carry = select(0u, 1u, sum < val);
+        }
+        for (var i = 8u; i < 16u; i++) {
+            current_nonce[i] = nonce_base[i];
+        }
+
+        // Resume the sponge from the precomputed midstate: absorb the low
+        // nonce half, pad, squeeze twice (3 permutations instead of 5).
+        var st: array<u64, 12>;
+        for (var i = 0u; i < 12u; i++) {
+            st[i] = mid[i];
+        }
+        for (var i = 0u; i < 8u; i++) {
+            st[i] = gf64_add(st[i], u64(bswap32(current_nonce[7u - i])));
+        }
+        permute64(&st);
+        st[0] = gf64_add(st[0], 1lu);
+        st[1] = gf64_add(st[1], 1lu);
+        permute64(&st);
+
+        // 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 candidates pay for the second squeeze, and
+        // byte-swapped hash words are produced on demand during the compare.
+        var first: array<u32, 8>;
+        for (var i = 0u; i < 4u; i++) {
+            let c = gf64_canon(st[i]);
+            first[2u * i] = u32(c & EPS64);
+            first[2u * i + 1u] = u32(c >> 32u);
+        }
+        var cmp = 0u;
+        for (var i = 0u; i < 8u; i++) {
+            let h = bswap32(first[i]);
+            let t = tgt[15u - i];
+            if (h != t) {
+                cmp = select(2u, 1u, h > t);
+                break;
+            }
+        }
+        if (cmp == 1u) {
+            continue;
+        }
+
+        var hash_le: array<u32, 16>;
+        for (var i = 0u; i < 8u; i++) {
+            hash_le[15u - i] = bswap32(first[i]);
+        }
+        permute64(&st);
+        for (var i = 0u; i < 4u; i++) {
+            let c = gf64_canon(st[i]);
+            hash_le[7u - 2u * i] = bswap32(u32(c & EPS64));
+            hash_le[6u - 2u * i] = bswap32(u32(c >> 32u));
+        }
+        var below = cmp == 2u;
+        if (!below) {
+            for (var i = 0u; i < 8u; i++) {
+                let h = hash_le[7u - i];
+                let t = tgt[7u - i];
+                if (h != t) {
+                    below = h < t;
+                    break;
+                }
+            }
+        }
+
+        if (below) {
+            if (atomicExchange(&results[0], 1u) == 0u) {
+                for (var i = 0u; i < 16u; i++) {
+                    atomicStore(&results[1u + i], current_nonce[i]);
+                    atomicStore(&results[17u + i], hash_le[i]);
+                }
+            }
+            return;
+        }
+    }
+}
+
+// ---------------------------------------------------------------------------
+// Compatibility layer: same API as mining.wgsl, backed by the u64 core above.
+// Only used by the component test harness; the mining kernel never calls it.
+// All outputs are canonical, matching the reference implementation.
+// ---------------------------------------------------------------------------
+
+struct GoldilocksField {
+    limb0: u32,
+    limb1: u32,
+}
+
+const INTERNAL_CONSTANTS: array<array<u32, 2>, 22> = array<array<u32, 2>, 22>(
+    array<u32, 2>(2018170979u, 2549578122u),
+    array<u32, 2>(794875120u, 3520249608u),
+    array<u32, 2>(2677723654u, 1772320679u),
+    array<u32, 2>(2743438884u, 2849007878u),
+    array<u32, 2>(518907317u, 693269760u),
+    array<u32, 2>(293328710u, 1484055617u),
+    array<u32, 2>(2834138828u, 2315799483u),
+    array<u32, 2>(1558078501u, 1039128420u),
+    array<u32, 2>(2266808631u, 966316006u),
+    array<u32, 2>(3393728842u, 1045622667u),
+    array<u32, 2>(2245828300u, 2521440415u),
+    array<u32, 2>(751064958u, 1070374632u),
+    array<u32, 2>(3490278765u, 2390340773u),
+    array<u32, 2>(3526960470u, 2224174634u),
+    array<u32, 2>(639988950u, 4000511088u),
+    array<u32, 2>(1839350858u, 504240201u),
+    array<u32, 2>(559852230u, 255489215u),
+    array<u32, 2>(2713771731u, 453385078u),
+    array<u32, 2>(1745082278u, 422331096u),
+    array<u32, 2>(2544763488u, 4141129721u),
+    array<u32, 2>(2700752774u, 1052996327u),
+    array<u32, 2>(4063512019u, 1429786100u)
+);
+
+const INITIAL_EXTERNAL_CONSTANTS: array<array<array<u32, 2>, 12>, 4> = array<array<array<u32, 2>, 12>, 4>(
+    array<array<u32, 2>, 12>(
+        array<u32, 2>(2539329031u, 3221415792u),
+        array<u32, 2>(4262746426u, 3164936845u),
+        array<u32, 2>(3883202553u, 1922272763u),
+        array<u32, 2>(3761386668u, 3841130025u),
+        array<u32, 2>(1411081289u, 3588274735u),
+        array<u32, 2>(4090250945u, 3962812520u),
+        array<u32, 2>(1100826458u, 1215155029u),
+        array<u32, 2>(1489773809u, 1813820067u),
+        array<u32, 2>(2585015995u, 3824356688u),
+        array<u32, 2>(2378857513u, 3651555078u),
+        array<u32, 2>(2864423342u, 3852156759u),
+        array<u32, 2>(1531416540u, 708695120u)
+    ),
+    array<array<u32, 2>, 12>(
+        array<u32, 2>(1987505445u, 2913073372u),
+        array<u32, 2>(1426707734u, 655469195u),
+        array<u32, 2>(3385403543u, 1256631504u),
+        array<u32, 2>(1381422714u, 1458257259u),
+        array<u32, 2>(2453402910u, 528129365u),
+        array<u32, 2>(964329320u, 905986685u),
+        array<u32, 2>(1534247888u, 3842469367u),
+        array<u32, 2>(744525997u, 4241857185u),
+        array<u32, 2>(1756723870u, 3448331916u),
+        array<u32, 2>(3610291774u, 1105166073u),
+        array<u32, 2>(2596181885u, 3997051784u),
+        array<u32, 2>(3199845381u, 3533420525u)
+    ),
+    array<array<u32, 2>, 12>(
+        array<u32, 2>(4127777666u, 2103183598u),
+        array<u32, 2>(2867888172u, 2309916828u),
+        array<u32, 2>(1831532055u, 3009056407u),
+        array<u32, 2>(2947502451u, 3675530062u),
+        array<u32, 2>(3565886616u, 2029012066u),
+        array<u32, 2>(3833391242u, 642945968u),
+        array<u32, 2>(1773785903u, 2577032347u),
+        array<u32, 2>(1770914259u, 1689297286u),
+        array<u32, 2>(3752758200u, 3993707216u),
+        array<u32, 2>(3389302766u, 1339375184u),
+        array<u32, 2>(2180141127u, 1466089441u),
+        array<u32, 2>(3199591357u, 4111832034u)
+    ),
+    array<array<u32, 2>, 12>(
+        array<u32, 2>(1625498743u, 509487959u),
+        array<u32, 2>(4188712685u, 1646551713u),
+        array<u32, 2>(3451003566u, 2854767422u),
+        array<u32, 2>(1412166652u, 1674110767u),
+        array<u32, 2>(3410212320u, 1000704202u),
+        array<u32, 2>(3381743837u, 602777331u),
+        array<u32, 2>(3131873882u, 2866003620u),
+        array<u32, 2>(2610174026u, 3923414377u),
+        array<u32, 2>(3644719692u, 3450945356u),
+        array<u32, 2>(1458984419u, 2418851081u),
+        array<u32, 2>(3344519983u, 1531855103u),
+        array<u32, 2>(2721413879u, 3732495392u)
+    )
+);
+
+const TERMINAL_EXTERNAL_CONSTANTS: array<array<array<u32, 2>, 12>, 4> = array<array<array<u32, 2>, 12>, 4>(
+    array<array<u32, 2>, 12>(
+        array<u32, 2>(3773451374u, 2071119161u),
+        array<u32, 2>(3805190518u, 340095962u),
+        array<u32, 2>(2402679944u, 2149591222u),
+        array<u32, 2>(743434178u, 1832305922u),
+        array<u32, 2>(2847530739u, 2718290175u),
+        array<u32, 2>(514243119u, 4142392203u),
+        array<u32, 2>(3844443492u, 888639642u),
+        array<u32, 2>(2008645578u, 2957397405u),
+        array<u32, 2>(3732799654u, 1692252629u),
+        array<u32, 2>(664231319u, 248567644u),
+        array<u32, 2>(287781771u, 482031345u),
+        array<u32, 2>(3486561978u, 1718871301u)
+    ),
+    array<array<u32, 2>, 12>(
+        array<u32, 2>(814165505u, 1616875560u),
+        array<u32, 2>(2123759183u, 1070984082u),
+        array<u32, 2>(2722916813u, 3893372341u),
+        array<u32, 2>(3726899022u, 4157656693u),
+        array<u32, 2>(2824360073u, 4086907574u),
+        array<u32, 2>(4155973110u, 1837140488u),
+        array<u32, 2>(2297731723u, 4169165669u),
+        array<u32, 2>(707924090u, 1474243980u),
+        array<u32, 2>(1298483757u, 384287239u),
+        array<u32, 2>(4243798069u, 557703745u),
+        array<u32, 2>(1510569718u, 2968696976u),
+        array<u32, 2>(3174388759u, 3638808363u)
+    ),
+    array<array<u32, 2>, 12>(
+        array<u32, 2>(2044277519u, 3835193622u),
+        array<u32, 2>(2743212840u, 1983595986u),
+        array<u32, 2>(3943309069u, 151568229u),
+        array<u32, 2>(424355060u, 1989274413u),
+        array<u32, 2>(867046322u, 239293714u),
+        array<u32, 2>(4230997871u, 2479068123u),
+        array<u32, 2>(1565052394u, 2566260552u),
+        array<u32, 2>(815274432u, 3822673712u),
+        array<u32, 2>(1051683535u, 519405993u),
+        array<u32, 2>(2687564964u, 186958263u),
+        array<u32, 2>(1450226471u, 1648586942u),
+        array<u32, 2>(1511122054u, 1595811937u)
+    ),
+    array<array<u32, 2>, 12>(
+        array<u32, 2>(2426274859u, 4261676319u),
+        array<u32, 2>(1207777283u, 1918363057u),
+        array<u32, 2>(3090099832u, 3870167883u),
+        array<u32, 2>(4101522705u, 1460796764u),
+        array<u32, 2>(201900220u, 4164567654u),
+        array<u32, 2>(2587682901u, 752404845u),
+        array<u32, 2>(2967564913u, 2100296475u),
+        array<u32, 2>(3404347409u, 2242778408u),
+        array<u32, 2>(3350048952u, 1386431957u),
+        array<u32, 2>(4093308564u, 1347177553u),
+        array<u32, 2>(2633812729u, 3169012324u),
+        array<u32, 2>(1727753673u, 3768793234u)
+    )
+);
+
+fn gf_pack(g: GoldilocksField) -> u64 {
+    return (u64(g.limb1) << 32u) | u64(g.limb0);
+}
+
+fn gf_unpack(v: u64) -> GoldilocksField {
+    let c = gf64_canon(v);
+    return GoldilocksField(u32(c & EPS64), u32(c >> 32u));
+}
+
+fn gf_from_limbs(l0: u32, l1: u32) -> GoldilocksField {
+    return GoldilocksField(l0, l1);
+}
+
+fn gf_zero() -> GoldilocksField {
+    return GoldilocksField(0u, 0u);
+}
+
+fn gf_one() -> GoldilocksField {
+    return GoldilocksField(1u, 0u);
+}
+
+fn gf_from_u32(val: u32) -> GoldilocksField {
+    return GoldilocksField(val, 0u);
+}
+
+fn gf_from_u64_parts(low: u32, high: u32) -> GoldilocksField {
+    return gf_unpack((u64(high) << 32u) | u64(low));
+}
+
+fn gf_from_const(val: array<u32, 2>) -> GoldilocksField {
+    return gf_from_u64_parts(val[0], val[1]);
+}
+
+fn gf_add(a: GoldilocksField, b: GoldilocksField) -> GoldilocksField {
+    return gf_unpack(gf64_add(gf_pack(a), gf_pack(b)));
+}
+
+fn gf_mul(a: GoldilocksField, b: GoldilocksField) -> GoldilocksField {
+    return gf_unpack(gf64_mul(gf_pack(a), gf_pack(b)));
+}
+
+fn sbox(x: GoldilocksField) -> GoldilocksField {
+    return gf_unpack(gf64_sbox(gf_pack(x)));
+}
+
+fn state_pack(state: ptr<function, array<GoldilocksField, 12>>, out: ptr<function, array<u64, 12>>) {
+    for (var i = 0u; i < 12u; i++) {
+        (*out)[i] = gf_pack((*state)[i]);
+    }
+}
+
+fn state_unpack(v: ptr<function, array<u64, 12>>, state: ptr<function, array<GoldilocksField, 12>>) {
+    for (var i = 0u; i < 12u; i++) {
+        (*state)[i] = gf_unpack((*v)[i]);
+    }
+}
+
+fn external_linear_layer(state: ptr<function, array<GoldilocksField, 12>>) {
+    var st: array<u64, 12>;
+    state_pack(state, &st);
+    ext_layer64(&st);
+    state_unpack(&st, state);
+}
+
+fn internal_linear_layer(state: ptr<function, array<GoldilocksField, 12>>) {
+    var st: array<u64, 12>;
+    state_pack(state, &st);
+    int_layer64(&st);
+    state_unpack(&st, state);
+}
+
+fn poseidon2_permute(state: ptr<function, array<GoldilocksField, 12>>) {
+    var st: array<u64, 12>;
+    state_pack(state, &st);
+    permute64(&st);
+    state_unpack(&st, state);
+}
+
+fn bytes_to_field_elements(input: array<u32, 24>) -> array<GoldilocksField, 25> {
+    var felts: array<GoldilocksField, 25>;
+    for (var i = 0u; i < 24u; i++) {
+        felts[i] = gf_from_u32(input[i]);
+    }
+    felts[24] = gf_one();
+    return felts;
+}
+
+fn field_elements_to_bytes(felts: array<GoldilocksField, 4>) -> array<u32, 8> {
+    var result: array<u32, 8>;
+    for (var i = 0u; i < 4u; i++) {
+        result[i * 2u] = felts[i].limb0;
+        result[i * 2u + 1u] = felts[i].limb1;
+    }
+    return result;
+}
+
+fn poseidon2_hash_squeeze_twice(input: array<u32, 24>) -> array<u32, 16> {
+    var st: array<u64, 12>;
+    for (var i = 0u; i < 12u; i++) {
+        st[i] = 0lu;
+    }
+    for (var chunk = 0u; chunk < 3u; chunk++) {
+        for (var i = 0u; i < 8u; i++) {
+            st[i] = gf64_add(st[i], u64(input[chunk * 8u + i]));
+        }
+        permute64(&st);
+    }
+    st[0] = gf64_add(st[0], 1lu);
+    st[1] = gf64_add(st[1], 1lu);
+    permute64(&st);
+
+    var result: array<u32, 16>;
+    for (var i = 0u; i < 4u; i++) {
+        let c = gf64_canon(st[i]);
+        result[2u * i] = u32(c & EPS64);
+        result[2u * i + 1u] = u32(c >> 32u);
+    }
+    permute64(&st);
+    for (var i = 0u; i < 4u; i++) {
+        let c = gf64_canon(st[i]);
+        result[8u + 2u * i] = u32(c & EPS64);
+        result[8u + 2u * i + 1u] = u32(c >> 32u);
+    }
+    return result;
+}
+
+fn hash_squeeze_twice(input: array<u32, 24>) -> array<u32, 16> {
+    return poseidon2_hash_squeeze_twice(input);
+}

crates/engine-gpu/src/kernels/mining_u64_apple.wgsl

@@ -1,7 +1,5 @@
-// Quantus Mining Shader - native-u64 variant (requires wgpu Features::SHADER_INT64)
-// Implements Poseidon2 hash over the Goldilocks field with plonky2-style lazy reduction:
-// values live in [0, 2^64) and are only canonicalized when squeezed out.
-// Must produce byte-identical results to mining.wgsl / qp-poseidon-core.
+// Native-u64 Poseidon2 mining kernel. Apple Metal path only.
+// Requires wgpu Features::SHADER_INT64. Bit-exact with mining.wgsl / pow_core.
 
 @group(0) @binding(0) var<storage, read_write> results: array<atomic<u32>>;
 // Sponge state after absorbing header + high nonce half (12 felts as LE u32 pairs),

crates/engine-gpu/src/kernels/mod.rs

@@ -0,0 +1,116 @@
+//! Poseidon2 mining kernels.
+//!
+//! Same `mining_main` bindings; must stay bit-exact with `pow_core`.
+//!
+//! - Apple Metal + `SHADER_INT64` → Apple Metal u64 (`mining_u64_apple.wgsl`)
+//! - other GPUs + `SHADER_INT64` → native u64 (`mining_u64.wgsl`)
+//! - no `SHADER_INT64` → 32-bit fallback (`mining.wgsl`)
+
+#[derive(Clone, Copy, Debug, Eq, PartialEq)]
+pub enum Kernel {
+    U32,
+    Default,
+    Apple,
+}
+
+impl Kernel {
+    pub const fn needs_int64(self) -> bool {
+        !matches!(self, Self::U32)
+    }
+
+    pub const fn id(self) -> &'static str {
+        match self {
+            Self::U32 => "u32",
+            Self::Default => "u64",
+            Self::Apple => "u64-apple",
+        }
+    }
+
+    pub const fn label(self) -> &'static str {
+        match self {
+            Self::U32 => "32-bit",
+            Self::Default => "native-u64",
+            Self::Apple => "native-u64 Apple Metal",
+        }
+    }
+
+    pub const fn source(self) -> &'static str {
+        match self {
+            Self::U32 => include_str!("mining.wgsl"),
+            Self::Default => include_str!("mining_u64.wgsl"),
+            Self::Apple => include_str!("mining_u64_apple.wgsl"),
+        }
+    }
+
+    pub fn for_adapter(adapter: &wgpu::Adapter) -> Self {
+        Self::for_adapter_info(&adapter.get_info(), adapter.features())
+    }
+
+    pub fn for_adapter_info(info: &wgpu::AdapterInfo, features: wgpu::Features) -> Self {
+        if !features.contains(wgpu::Features::SHADER_INT64) {
+            return Self::U32;
+        }
+        if info.backend == wgpu::Backend::Metal {
+            Self::Apple
+        } else {
+            Self::Default
+        }
+    }
+
+    pub const fn all() -> &'static [Self] {
+        &[Self::U32, Self::Default, Self::Apple]
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    fn info(backend: wgpu::Backend) -> wgpu::AdapterInfo {
+        wgpu::AdapterInfo {
+            name: "test".into(),
+            vendor: 0,
+            device: 0,
+            device_type: wgpu::DeviceType::DiscreteGpu,
+            driver: String::new(),
+            driver_info: String::new(),
+            backend,
+        }
+    }
+
+    #[test]
+    fn metal_with_int64_selects_apple() {
+        assert_eq!(
+            Kernel::for_adapter_info(&info(wgpu::Backend::Metal), wgpu::Features::SHADER_INT64),
+            Kernel::Apple
+        );
+    }
+
+    #[test]
+    fn vulkan_with_int64_selects_default() {
+        assert_eq!(
+            Kernel::for_adapter_info(&info(wgpu::Backend::Vulkan), wgpu::Features::SHADER_INT64),
+            Kernel::Default
+        );
+    }
+
+    #[test]
+    fn dx12_with_int64_selects_default() {
+        assert_eq!(
+            Kernel::for_adapter_info(&info(wgpu::Backend::Dx12), wgpu::Features::SHADER_INT64),
+            Kernel::Default
+        );
+    }
+
+    #[test]
+    fn no_int64_falls_back_to_u32() {
+        assert_eq!(
+            Kernel::for_adapter_info(&info(wgpu::Backend::Metal), wgpu::Features::empty()),
+            Kernel::U32
+        );
+        assert_eq!(
+            Kernel::for_adapter_info(&info(wgpu::Backend::Vulkan), wgpu::Features::empty()),
+            Kernel::U32
+        );
+    }
+}

crates/engine-gpu/src/lib.rs

@@ -2,10 +2,13 @@
 #![deny(unsafe_code)]
 
 mod gpu_tiers;
+mod kernels;
 
 pub mod end_to_end_tests;
 pub mod tests;
 
+pub use kernels::Kernel;
+
 use engine_cpu::{CancelCheck, Candidate, EngineStatus, FoundOrigin, MinerEngine, Range};
 use pow_core::{format_hashrate, format_u512, JobContext};
 use primitive_types::U512;
@@ -17,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,
@@ -153,10 +158,10 @@ impl GpuContext {
 /// Create the mining shader module without naga's runtime bounds checks and
 /// loop bounding (~9% faster kernels).
 ///
-/// SAFETY: the sources are the static mining shaders compiled into this binary;
+/// SAFETY: the sources are the static mining kernels compiled into this binary;
 /// every buffer access is a constant-bounded loop index into fixed-size
 /// bindings the engine itself allocates, and all loops have static bounds
-/// (verified by the component test suites against both shader variants).
+/// (verified by the component test suites against every kernel path).
 #[allow(unsafe_code)]
 fn create_trusted_shader(device: &wgpu::Device, shader_source: &str) -> wgpu::ShaderModule {
     unsafe {
@@ -357,15 +362,13 @@ impl GpuEngine {
             );
             log::debug!(target: "gpu_engine", "Adapter {i} raw info: {info:?}");
 
-            // Prefer the native-u64 shader where supported (Apple/NVIDIA/modern AMD):
-            // Goldilocks arithmetic on u64 is far cheaper than 32-bit limb emulation.
-            let use_u64 = adapter.features().contains(wgpu::Features::SHADER_INT64);
+            let kernel = Kernel::for_adapter(&adapter);
 
             // Try to initialize this adapter with a proper timeout.
             // If the driver hangs, we'll skip this adapter after the timeout.
             let device_future = adapter.request_device(&wgpu::DeviceDescriptor {
                 label: Some("Mining Device"),
-                required_features: if use_u64 {
+                required_features: if kernel.needs_int64() {
                     wgpu::Features::SHADER_INT64
                 } else {
                     wgpu::Features::empty()
@@ -412,18 +415,14 @@ impl GpuEngine {
             // Shader and pipeline creation are synchronous - can't timeout, but usually fast
             let pipeline_start = std::time::Instant::now();
 
-            let shader_source = if use_u64 {
-                include_str!("mining_u64.wgsl")
-            } else {
-                include_str!("mining.wgsl")
-            };
             log::info!(
                 target: "gpu_engine",
-                "GPU device {i} ({}) using {} shader",
+                "GPU device {i} ({}) using {} [{}]",
                 info.name,
-                if use_u64 { "native-u64" } else { "32-bit" }
+                kernel.label(),
+                kernel.id()
             );
-            let shader = create_trusted_shader(&device, shader_source);
+            let shader = create_trusted_shader(&device, kernel.source());
 
             let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
                 label: Some("Mining Pipeline"),
@@ -446,6 +445,7 @@ impl GpuEngine {
 
             initialized.push(InitializedGpu {
                 context: Arc::new(GpuContext {
+                    metrics: metrics::DeviceMetrics::new(i, kernel.id()),
                     device,
                     queue,
                     pipeline,
@@ -618,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",
@@ -632,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!(
@@ -699,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);
 
@@ -796,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;
@@ -868,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(..);
@@ -914,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();
@@ -942,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,
@@ -951,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/engine-gpu/src/main.rs

@@ -13,37 +13,28 @@ async fn main() -> Result<(), Box<dyn std::error::Error>> {
         .expect("no GPU adapter");
 
     let mut failures = 0usize;
+    let has_int64 = adapter.features().contains(wgpu::Features::SHADER_INT64);
 
-    {
-        let (device, queue) = adapter
-            .request_device(&wgpu::DeviceDescriptor::default())
-            .await?;
-        failures += run_suite(
-            &device,
-            &queue,
-            include_str!("mining.wgsl"),
-            "32-bit (mining.wgsl)",
-        )
-        .await;
-    }
-
-    if adapter.features().contains(wgpu::Features::SHADER_INT64) {
+    for kernel in engine_gpu::Kernel::all() {
+        if kernel.needs_int64() && !has_int64 {
+            println!(
+                "\nSHADER_INT64 not supported on this adapter; skipping {}",
+                kernel.label()
+            );
+            continue;
+        }
         let (device, queue) = adapter
             .request_device(&wgpu::DeviceDescriptor {
-                label: Some("u64 Test Device"),
-                required_features: wgpu::Features::SHADER_INT64,
+                label: Some(kernel.label()),
+                required_features: if kernel.needs_int64() {
+                    wgpu::Features::SHADER_INT64
+                } else {
+                    wgpu::Features::empty()
+                },
                 ..Default::default()
             })
             .await?;
-        failures += run_suite(
-            &device,
-            &queue,
-            include_str!("mining_u64.wgsl"),
-            "native-u64 (mining_u64.wgsl)",
-        )
-        .await;
-    } else {
-        println!("\nSHADER_INT64 not supported on this adapter; skipping mining_u64.wgsl suite");
+        failures += run_suite(&device, &queue, kernel.source(), kernel.label()).await;
     }
 
     if failures > 0 {

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,8 @@
+# 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}}
+# Kernel entry in the CUDA engine: unrolled (mining_main) or loop (mining_loop), quantus/miner#29.
+MINER_CUDA_KERNEL={{MINER_CUDA_KERNEL}}

deploy/quantus-miner-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

rust-toolchain

@@ -1,4 +1,4 @@
 [toolchain]
-channel = "stable"
+channel = "1.93.0"
 components = ["clippy", "rustfmt"]
 profile = "minimal"