blackbeard/miner
1
0
Fork 0
Code Issues 2 Pull Requests Actions Packages Projects Releases 2 Wiki Activity

Compare commits

base: blackbeard/miner:v3.3.0
Branches Tags
blackbeard/miner:main blackbeard/miner:perf/minblocks-knob blackbeard/miner:perf/sm120-spills blackbeard/miner:origin-main
blackbeard/miner:v4.0.2-lair.2 blackbeard/miner:v4.0.2-lair.1 blackbeard/miner:v4.0.2 blackbeard/miner:v4.0.1 blackbeard/miner:v4.0.0 blackbeard/miner:v3.3.1 blackbeard/miner:v3.3.0 blackbeard/miner:v3.2.0 blackbeard/miner:v3.1.0 blackbeard/miner:v3.0.1 blackbeard/miner:v3.0.0 blackbeard/miner:v2.1.2 blackbeard/miner:v2.1.1 blackbeard/miner:v2.1.0 blackbeard/miner:v2.0.2 blackbeard/miner:v2.0.1 blackbeard/miner:v2.0.0 blackbeard/miner:v1.0.0 blackbeard/miner:v0.3.0 blackbeard/miner:v0.2.1
...
compare: blackbeard/miner:v4.0.2
Branches Tags
blackbeard/miner:main blackbeard/miner:perf/minblocks-knob blackbeard/miner:perf/sm120-spills blackbeard/miner:origin-main
blackbeard/miner:v4.0.2-lair.2 blackbeard/miner:v4.0.2-lair.1 blackbeard/miner:v4.0.2 blackbeard/miner:v4.0.1 blackbeard/miner:v4.0.0 blackbeard/miner:v3.3.1 blackbeard/miner:v3.3.0 blackbeard/miner:v3.2.0 blackbeard/miner:v3.1.0 blackbeard/miner:v3.0.1 blackbeard/miner:v3.0.0 blackbeard/miner:v2.1.2 blackbeard/miner:v2.1.1 blackbeard/miner:v2.1.0 blackbeard/miner:v2.0.2 blackbeard/miner:v2.0.1 blackbeard/miner:v2.0.0 blackbeard/miner:v1.0.0 blackbeard/miner:v0.3.0 blackbeard/miner:v0.2.1
23 Commits
v3.3.0 ... v4.0.2
Author SHA1 Message Date
Nikolaus Hegerandn13 cb6deb9139 bump version to v4.0.2 (#93)
CI / 🏁 Fast Checks (Format) (push) Has been cancelled
Details
CI / 🛠️ Build & Test (push) Has been cancelled
Details
CI / 🤖 Analysis (Clippy & Doc) (push) Has been cancelled
Details
CI / 🏃 Benchmark (push) Has been cancelled
Details 
Co-authored-by: n13 <n13@users.noreply.github.com>
 2026-09-01 23:25:24 +08:00
Nikolaus Heger 09323d33e4 Different kernels for Apple Metal vs Nvidia and everything else (#92) 
* Select Apple Metal 4.0.1 kernel; keep 4.0.0 u64 elsewhere

NVIDIA was ~37% slower on the 4.0.1 Metal-tuned kernel. Split
Poseidon2 kernels and pick by backend:

- Metal + SHADER_INT64: v4.0.1 u64
- other SHADER_INT64: v4.0.0 u64
- no SHADER_INT64: 32-bit fallback (same as 4.0.0)

CI uploads the Linux release binary as an artifact so we can
bench NVIDIA without building on the GPU box.

* Drop version numbers from GPU kernel labels

Log native-u64 vs native-u64 Apple Metal without tying
the strings to release tags.
 2026-09-01 23:21:43 +08:00
Nikolaus Heger ceff470a5f ci: pin the stable rust toolchain (#91) 
CI resolved the floating stable channel, so Rust 1.98 shipped the
chunks_exact_to_as_chunks clippy lint and failed checks that passed on
an older local stable. Pin the exact version, matching how chain pins
its toolchains.
 2026-09-01 19:03:54 +08:00
Nikolaus Hegerandn13 41ec132139 bump version to v4.0.1 (#89) 
Co-authored-by: n13 <n13@users.noreply.github.com>
 2026-09-01 19:03:05 +08:00
Nikolaus Heger ea5e5e742e Speed up the u64 mining kernel ~1.8x on Apple Metal (#87) 
Apple M5 Pro, CLI benchmark: 24.0 -> 43.2 MH/s. Kernel output is
unchanged and bit-exact against the CPU reference.

Field arithmetic:
- Deferred-carry additions (Acc): the external layer and the internal
  row sum accumulate unreduced with a carry counter and fold once per
  output instead of two epsilon corrections per add.
- gf64_mul_add folds the row sum into the 128-bit product before a
  single plain reduction.
- mul_wide / gf64_sqr assemble the 128-bit product from 32-bit-valued
  partial sums with no carry compares.

Kernel structure (code size dominates on Apple's compiler; every
unrolled variant measured slower):
- One loop drives all 30 rounds so each layer is emitted once, with
  round constants added by the linear layer preceding each S-box
  (RC_EXT table, RC_INTERNAL padded with a trailing zero).
- One S-box site with a runtime lane count (12 external, 1 internal).
- mining_main runs pad / first squeeze / second squeeze through a
  single inlined permute64 call in a phase loop.

Measured and rejected: uniform buffers and unhoisted loads, workgroup
sizes 64/128/512, 8-add MDS, one-compare reduce, fully unrolled
layers, two nonces per thread.
 2026-08-31 22:50:27 +08:00
I Dewa Gede Bisma Mahendra 696eb51125 feat: add dependency cooldown (#88) 2026-08-28 14:11:24 +08:00
Nikolaus Heger 8d0c9f8023 docs: update systemd example for the serve CLI with required node auth (#86) 
* docs: update systemd example for the serve CLI with required node auth

The unit and readme still described the old HTTP-API miner (MINER_PORT,
MINER_ENGINE, no subcommand). Update ExecStart to 'serve' and document the
now-required MINER_AUTH_TOKEN_FILE / MINER_TLS_CERT_SHA256_FILE, including
how to copy the node's miner-auth-token and miner-tls-cert-sha256 files
past ProtectHome=true, and refresh the env var reference to the current
CLI.

* docs: fix env-file example and CPU-worker claims per review

- Move inline comments off the env assignments: systemd EnvironmentFile=
  keeps inline '# ...' text as part of the value, so the copied example fed
  Clap unparseable values and the unit crash-looped under Restart=always.
- Describe the real worker behavior: CPUs are counted via the process
  affinity mask (num_cpus); unset MINER_CPU_WORKERS auto-detects ~50%, and
  an explicit value is used as-is — there is no clamp or warning.
- Drop the 'logs the detected cpuset mask' claim from the unit and override
  comments; no such logging exists in the miner.

* docs: metrics exporter is always on; purge-chain does not rotate credentials

- The Prometheus exporter starts unconditionally and binds plaintext HTTP on
  0.0.0.0:9900 by default; MINER_METRICS_PORT only changes the port. Say so
  everywhere instead of 'when metrics are enabled', and add a hardening note
  to firewall the port.
- Rephrase the credential-rotation pitfall: purge-chain removes only the
  database, so the node reloads the same miner-auth-token and TLS cert;
  rotation only happens when the base path changes or the files themselves
  are deleted.
 2026-08-14 11:46:35 +08:00
illuzenandn13 37f7837f76 bump version to v4.0.0 (#85) 
Co-authored-by: n13 <n13@users.noreply.github.com>
 2026-08-13 19:34:16 +08:00
illuzenandCursor 2c8d530de4 Send miner auth token in Ready handshake. (#78) 
* Send miner auth token in Ready handshake.

Require --auth-token or --auth-token-file when connecting to a node so the
miner matches the node's shared-secret QUIC auth.

Co-authored-by: Cursor <cursoragent@cursor.com>

* Pin the node's miner TLS certificate by SHA-256 fingerprint.

Require --tls-cert-sha256 or --tls-cert-sha256-file so miners reject
MITM certs instead of accepting any self-signed server certificate.

Co-authored-by: Cursor <cursoragent@cursor.com>

* fmt

* Docs: prefer reading miner auth token from file, not node logs.

Co-authored-by: Cursor <cursoragent@cursor.com>

* Version the miner ALPN as quantus-miner/2 to match the node.

A node/miner protocol mismatch now fails cleanly at the TLS handshake with
"no application protocol" instead of an opaque auth error.

Co-authored-by: Cursor <cursoragent@cursor.com>

* point at git

* Depend on published quantus-miner-api 0.3.0 from crates.io.

The temporary git branch pin is no longer needed now that 0.3.0 is published.

Co-authored-by: Cursor <cursoragent@cursor.com>

* Fail closed on permanent auth/TLS misconfig instead of reconnect looping.

Validate the fingerprint and Ready frame size at miner/pool startup before
metrics, workers, or HTTP come up. Treat node auth rejection and pin mismatch
as PermanentConnectError (no retry), and only reset reconnect backoff after
the first NewJob proves auth succeeded.

Co-authored-by: Cursor <cursoragent@cursor.com>

* Point miner auth docs and examples at miner-auth-token / TLS pin files.

The node never logs the auth token; CLI/pool help and every serve/Docker
example now require the chain config files (or env vars) instead of implying
credentials can be copied from logs.

Co-authored-by: Cursor <cursoragent@cursor.com>

* fmt

* Remove the miner Dockerfile and Docker packaging docs/CI.

Container GPU access is a poor fit for this miner, and the Docker examples
were unusable with SocketAddr (hostnames rejected). Run the native binary
instead.

Co-authored-by: Cursor <cursoragent@cursor.com>

* Shut down the pool on permanent upstream auth failure; keep reconnecting otherwise.

Supervise run_node_client alongside HTTP so a bad token/pin ends the process
instead of leaving a live API on a dead upstream. Node restarts still reconnect
with backoff; the current job is cleared while disconnected so captchas are not
issued against a stale header.

Co-authored-by: Cursor <cursoragent@cursor.com>

* Replace obsolete EXTERNAL_MINER_PROTOCOL.md with a pointer to MINING.md.

The local copy still described the pre-auth protocol (no Ready token, no pin,
old ALPN, 16 MB frames). Canonical docs live in quantus-miner-api and the
node's MINING.md.

Co-authored-by: Cursor <cursoragent@cursor.com>

* Delete EXTERNAL_MINER_PROTOCOL.md; README already points at MINING.md.

Co-authored-by: Cursor <cursoragent@cursor.com>

---------

Co-authored-by: Cursor <cursoragent@cursor.com>
 2026-08-13 09:50:03 +08:00
Nikolaus Heger 2b966e780a Revert "Fuzz GPU/CPU parity at scale and CPU-verify every seal before submission" 
This reverts commit b0cfc3005e.
 2026-08-12 15:41:38 +08:00
Nikolaus HegerandClaude Fable 5 b0cfc3005e Fuzz GPU/CPU parity at scale and CPU-verify every seal before submission 
Investigating reported seal submission failures. Root cause analysis:
the miner's job-id staleness chain (worker check, quic-loop filter,
node-side job_id compare) cannot mislabel a candidate — internal and
node job ids are updated atomically together. The remaining benign race
is node-side: a block-template rebuild (new pre_hash, same parent) does
not bump the node's job counter, so an in-flight seal mined on the old
template passes the job-id check and fails seal verification, logging
"Failed to submit seal". Fixing that belongs in the chain repo (bump the
job counter on template rebuild).

Miner-side hardening in this commit:

- gpu_cpu_parity is now a two-phase fuzzer. Phase 1 dumps FULL 512-bit
  hashes for randomized (header, nonce) pairs via the real
  midstate-resume datapath and verifies each against the canonical CPU
  implementation, multi-threaded (defaults to 1M hashes, ~5s; 10M runs
  in ~15s). Phase 2 fuzzes whole mining jobs with five profiles: random,
  carry-edge starts (saturated low limbs), 2^256-boundary crossings,
  impossible difficulty with exact hash-count assertion, and CPU-known
  solutions to catch false negatives. Seeded and reproducible; also
  asserts the submitted work bytes match the nonce.
  Verified clean: 10M bulk hashes + ~1,500 seals across 6 seeds.

- miner-service re-verifies every candidate on the CPU (one hash) before
  sending JobResult; an engine-produced invalid seal is never submitted
  and is logged loudly instead.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
 2026-08-12 14:24:05 +08:00
Nikolaus Heger 9b16df1dff Merge pull request #81 from Quantus-Network/gpu-poseidon2-optimization 
GPU mining 4.1x faster on Apple Silicon: native-u64 Poseidon2, midstate precompute, lazy second squeeze
 2026-08-12 12:10:58 +08:00
Nikolaus Heger 7e449c0a5f Benchmark: honor --gpu-batch-size/--cpu-batch-size and advance nonce ranges 
The benchmark fed each worker a fixed range (10k CPU / 1M GPU nonces),
capping GPU dispatches at 1M regardless of --gpu-batch-size, and
re-searched the same range on every iteration. Worker ranges now step
through the nonce space and are sized from the batch flags, floored at
the old widths so tiny flags keep engine-internal cancellation batching
instead of turning per-call harness overhead into the measured quantity.
Batch-size flags are validated non-zero at parse time (serve and
benchmark) and echoed in the benchmark banner.
 2026-08-12 12:04:44 +08:00
illuzen 22d87d4772 fix benches 2026-08-12 11:16:26 +08:00
Nikolaus HegerandClaude Fable 5 8aa81ec853 Disable naga runtime shader checks for the trusted mining shaders (+8-9%) 
Exploring a fully native Metal port of the mining kernel showed it runs
~16% faster than the wgpu path at equal batch size. Decomposing that gap:
~9% is naga's injected runtime bounds checks and forced loop bounding,
~1.5% is mulhi()-based multiplication (inexpressible in WGSL), <1% is
the constant address space; the rest is binding/codegen residue.

The dominant term is recoverable inside wgpu on every backend:
the engine now builds its pipelines with create_shader_module_trusted +
ShaderRuntimeChecks::unchecked(). This is sound for our shaders: sources
are compiled into the binary, every buffer access is a constant-bounded
loop index into fixed-size bindings the engine allocates itself, all
loops have static bounds, and both variants remain covered by the
dual-shader component suite and the CPU parity example.

Apple M4: 11.39 -> 12.34 MH/s at the default 1M batch (+8.3%), 11.74 ->
12.73 MH/s at 8M (+8.4%). Cumulative vs main: 4.4x.

New examples/trusted_hashrate.rs measures checked vs trusted on the same
shader and CPU-verifies results.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
 2026-08-12 00:55:15 +08:00
Nikolaus HegerandClaude Fable 5 58615ed3fa Adopt PR #80 micro-optimizations that measure positive on the u64 kernel 
Evaluated all ideas from #80 against the native-u64 kernel. Its three
main wins (lazy field arithmetic, precomputed sponge state, conditional
second squeeze) were already present in stronger form; two micro-opts
carried over and measured positive at large batches (+1.3% at 16M,
noise-level at the 1M default):

- Hoist midstate/target/start-nonce storage reads out of the nonce loop.
- Produce byte-swapped hash words on demand during the target compare,
  so the reject path skips building the swapped hash entirely.

Forcing nonces_per_thread=2 was also tested and regressed; one thread
per nonce stays.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
 2026-08-12 00:01:39 +08:00
Nikolaus HegerandClaude Fable 5 d55b970ab6 Optimize GPU Poseidon2 mining 4.1x on Apple Silicon (u64 shader, midstate, lazy squeeze) 
Three stacked optimizations, each verified bit-exact against the CPU
reference (qp-poseidon-core) by the dual-shader component suite and a
new GPU/CPU parity example:

- mining_u64.wgsl: native 64-bit Goldilocks arithmetic with plonky2-style
  lazy reduction, used automatically when the adapter supports
  SHADER_INT64 (all Apple Silicon, NVIDIA, modern AMD). mining.wgsl
  remains the 32-bit fallback. (+50%)
- Midstate precompute: the first two of five sponge permutations absorb
  the header and high nonce half, both constant per batch; they are now
  computed once on the CPU (pow_core::mining_midstate) and resumed on
  the GPU. Batches are clamped so nonces never carry past 2^256. (+77%)
- Lazy second squeeze: the first squeeze yields the most significant 256
  bits of the hash, which decide hash-vs-target unless exactly equal to
  the target's high half, so the common reject path skips the final
  permutation. (+56%)

Apple M4 throughput: 2.81 -> 11.4 MH/s at the default 1M batch
(criterion large_range_1m/gpu: 355.8ms -> 86.3ms, -75.7%).

Also:
- Fix criterion GPU bench crash: thread-local worker resources are now
  tagged with an engine id so multiple GpuEngines per process never mix
  devices.
- Fix stale end-to-end test harness (6-binding layout) and run the full
  component suite against both shader variants.
- New examples: hashrate, gpu_cpu_parity (25 jobs incl. 2^256-boundary
  crossing), gpu_features.
- Add Apple M5 family GPU tiers.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
 2026-08-11 23:40:12 +08:00
Nikolaus Heger 0135e65278 Merge pull request #76 from Quantus-Network/add-security-policy 
Add security policy
 2026-08-11 22:08:43 +08:00
Nikolaus Heger 2eb0341a70 chore: bump qp-poseidon-core to 3.1.0 (#77) 
Update workspace pin from 2.1.0 to 3.1.0 for sponge wipe hardening.
(qpow-math from the pinned chain tag may still pull 2.1.0 transitively.)
 2026-08-08 15:47:42 +08:00
Nikolaus HegerandClaude Fable 5 0c61e3da34 Add security policy 
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
 2026-08-07 12:11:38 +08:00
illuzenandCursor 71781a0a80 Add captcha share pool service and browser WASM solver (#74) 
* Add captcha share pool service and browser WASM solver

pool-service connects to a node over the external-miner QUIC protocol and
turns browser captcha solves into real mining shares: sessions with disjoint
nonce ranges over the current header, single-use share tokens, a
reCAPTCHA-shaped /siteverify endpoint, and upstream block submission when a
share meets full network difficulty. solver-wasm is a raw C-ABI wasm32 build
of the Poseidon2 nonce grinder (no wasm-bindgen) for the embeddable widget.

Co-authored-by: Cursor <cursoragent@cursor.com>

* fmt

* Re-queue block solution when upstream write fails

A network-difficulty share taken off the solution channel was permanently
lost if the JobResult write to the node failed: the connection loop broke to
reconnect but the block was no longer anywhere. Park the block in a pending
slot that survives reconnection and is retried before any other work.

Co-authored-by: Cursor <cursoragent@cursor.com>

* fmt

* Address PR review: capacity limits, loud block drops, shared QUIC transport, constant-time secret compare

- Cap live sessions and tokens (configurable via --max-sessions/--max-tokens,
  default 100k each); /api/session returns 503 at_capacity so a request flood
  can no longer grow the maps without bound.
- A block solution that fails to queue upstream is now logged at error level,
  and blocks_found/block_found only report solutions actually queued.
- Extract the node QUIC connect + insecure-verifier pattern into a shared
  quic-transport crate used by both miner-service and pool-service; derive
  the network target via pow_core::JobContext instead of re-deriving it.
- Compare the site secret in constant time (constant_time_eq).
- Note single-tenant token semantics on site_secret for the host-registration
  follow-up.

Co-authored-by: Cursor <cursoragent@cursor.com>

* Add per-IP rate limiting on /api/session issuance

The global session cap stops unbounded map growth, but a single client could
still churn entries at the TTL boundary. Limit each IP to 60 issuances per
minute (configurable via --sessions-per-ip-per-min); excess requests get 429
rate_limited. Stale IP windows are GC'd alongside sessions/tokens.

Co-authored-by: Cursor <cursoragent@cursor.com>

* clippy

---------

Co-authored-by: Cursor <cursoragent@cursor.com>
 2026-07-07 17:52:38 +08:00
illuzen 15edf4a806 bump version to v3.3.1 (#73) 2026-06-30 16:34:53 +08:00
illuzen a47db3d1a2 Fix windows again (#72) 
* handle lost device correctly

* Update lib.rs

* allow-integrated flag otherwise exclude

* DeviceLost is a status

* Integrated GPU skipped when discrete init fails

* nits
 2026-06-30 16:30:11 +08:00
50 changed files with 4709 additions and 988 deletions
Expand all files Collapse all files
.dockerignore
-42
View File
@@ -1,42 +0,0 @@
# Git
.git
.gitignore
.github
# Build artifacts
target/
*.rs.bk
*.pdb
# IDE
.vscode
.idea
*.swp
*.swo
*~
# Documentation
docs/
*.md
!README.md
# Examples and tests (not needed for runtime)
examples/
# Other
.DS_Store
Thumbs.db
# CI/CD
.github/
# License
LICENSE
# Dockerfiles themselves
Dockerfile*
.dockerignore
# Cargo cache
.cargo/
.github/workflows/ci.yml
+5
View File
@@ -86,6 +86,11 @@ jobs:
- name: build benchmark binary
run: cargo build -p miner-cli --release
timeout-minutes: 60
- name: upload linux release binary
uses: actions/upload-artifact@v4
with:
name: quantus-miner-linux-x86_64
path: target/release/quantus-miner
- name: run cpu benchmark
run: ./target/release/quantus-miner benchmark --cpu-workers 2 --duration 5
timeout-minutes: 10
.github/workflows/dependency-cooldown-audit.yml
+13
View File
@@ -0,0 +1,13 @@
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
+14
View File
@@ -0,0 +1,14 @@
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/docker-image.yml
-139
View File
@@ -1,139 +0,0 @@
name: Docker Image
on:
workflow_dispatch:
inputs:
specific_version:
description: "Optional: Specify a full version (e.g., v0.3.0 or v0.3.1-beta.1) to build. If empty, uses the latest release tag. MUST start with 'v'."
required: false
default: ""
use_main_dockerfile:
description: "Optional: Use the main Dockerfile instead of the tag version-specific Dockerfile. If empty, uses the tag version-specific Dockerfile."
required: false
default: false
type: boolean
env:
CARGO_TERM_COLOR: always
jobs:
determine_version:
runs-on: ubuntu-latest
outputs:
version_with_v: ${{ steps.version_info.outputs.version_with_v }}
is_valid_version: ${{ steps.version_info.outputs.is_valid_version }}
steps:
- name: Checkout code
uses: actions/checkout@v4
with:
fetch-depth: 0
- name: Get version information
id: version_info
run: |
set -ex
TARGET_VERSION=""
SPECIFIC_VERSION="${{ github.event.inputs.specific_version }}"
if [[ -n "$SPECIFIC_VERSION" ]]; then
if [[ ! "$SPECIFIC_VERSION" =~ ^v[0-9]+\.[0-9]+\.[0-9]+(-[a-zA-Z0-9.-]+)?$ ]]; then
echo "::error::Specified version '$SPECIFIC_VERSION' is not a valid format. It must start with 'v' (e.g., v0.3.0 or v0.3.1-beta.1)."
echo "is_valid_version=false" >> $GITHUB_OUTPUT
exit 1
fi
TARGET_VERSION="$SPECIFIC_VERSION"
echo "Using specified version: $TARGET_VERSION"
else
echo "No specific version provided, determining latest release tag..."
TARGET_VERSION=$(git tag --list 'v*' --sort=-v:refname | head -n 1)
if [[ -z "$TARGET_VERSION" ]]; then
echo "::error::No version tags starting with 'v' (e.g., vX.Y.Z) found in the repository."
echo "is_valid_version=false" >> $GITHUB_OUTPUT
exit 1
fi
echo "Latest release version found: $TARGET_VERSION"
fi
echo "version_with_v=$TARGET_VERSION" >> $GITHUB_OUTPUT
echo "is_valid_version=true" >> $GITHUB_OUTPUT
build_and_publish_image:
needs: determine_version
if: needs.determine_version.outputs.is_valid_version == 'true'
runs-on: ubuntu-latest
permissions:
contents: read
packages: write
env:
GHCR_IMAGE_PATH: ghcr.io/quantus-network/quantus-miner
TARGET_VERSION_WITH_V: ${{ needs.determine_version.outputs.version_with_v }}
steps:
- name: Set up QEMU
uses: docker/setup-qemu-action@v3
- name: Set up Docker Buildx
uses: docker/setup-buildx-action@v3
- name: Log in to GitHub Container Registry
uses: docker/login-action@v3
with:
registry: ghcr.io
username: ${{ github.actor }}
password: ${{ secrets.GITHUB_TOKEN }}
- name: Check if Docker image already exists and fail if so
id: check_image
run: |
set -ex
IMAGE_TO_CHECK="${{ env.GHCR_IMAGE_PATH }}:${{ env.TARGET_VERSION_WITH_V }}"
echo "Checking for image: $IMAGE_TO_CHECK"
if docker manifest inspect "$IMAGE_TO_CHECK" > /dev/null 2>&1; then
echo "::error::Image $IMAGE_TO_CHECK already exists. Aborting."
exit 1
else
echo "Image $IMAGE_TO_CHECK does not exist. Proceeding with build."
fi
- name: Checkout code at tag version
uses: actions/checkout@v4
with:
ref: ${{ env.TARGET_VERSION_WITH_V }}
- name: Checkout main branch for Dockerfile (if needed)
if: github.event.inputs.use_main_dockerfile == 'true'
uses: actions/checkout@v4
with:
ref: main
sparse-checkout: |
Dockerfile
.dockerignore
sparse-checkout-cone-mode: false
path: dockerfile-source
- name: Copy Dockerfile from main (if needed)
if: github.event.inputs.use_main_dockerfile == 'true'
run: |
cp dockerfile-source/Dockerfile ./
[ -f dockerfile-source/.dockerignore ] && cp dockerfile-source/.dockerignore ./ || true
- name: Build and push Docker image
uses: docker/build-push-action@v6
with:
context: .
file: ./Dockerfile
push: true
platforms: linux/amd64,linux/arm64
provenance: false
tags: |
${{ env.GHCR_IMAGE_PATH }}:${{ env.TARGET_VERSION_WITH_V }}
${{ env.GHCR_IMAGE_PATH }}:latest
cache-from: type=gha
cache-to: type=gha,mode=max
- name: Print completion message
run: |
echo "Successfully built and pushed ${{ env.GHCR_IMAGE_PATH }}:${{ env.TARGET_VERSION_WITH_V }}"
echo "Image is also tagged as latest: ${{ env.GHCR_IMAGE_PATH }}:latest"
Cargo.lock
Generated
+84 -15
View File
@@ -377,6 +377,12 @@ dependencies = [
"unicode-xid",
]
[[package]]
name = "constant_time_eq"
version = "0.3.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "7c74b8349d32d297c9134b8c88677813a227df8f779daa29bfc29c183fe3dca6"
[[package]]
name = "core-foundation"
version = "0.9.4"
@@ -571,7 +577,7 @@ dependencies = [
[[package]]
name = "engine-cpu"
version = "3.3.0"
version = "4.0.2"
dependencies = [
"criterion",
"hex",
@@ -582,7 +588,7 @@ dependencies = [
[[package]]
name = "engine-gpu"
version = "3.3.0"
version = "4.0.2"
dependencies = [
"bytemuck",
"criterion",
@@ -596,7 +602,7 @@ dependencies = [
"qp-plonky2",
"qp-plonky2-field",
"qp-poseidon-constants",
"qp-poseidon-core",
"qp-poseidon-core 3.1.0",
"rand 0.9.2",
"rand_chacha 0.9.0",
"regex",
@@ -630,7 +636,7 @@ source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "39cab71617ae0d63f51a36d69f866391735b51691dbda63cf6f96d042b63efeb"
dependencies = [
"libc",
"windows-sys 0.61.2",
"windows-sys 0.52.0",
]
[[package]]
@@ -1270,7 +1276,7 @@ checksum = "3640c1c38b8e4e43584d8df18be5fc6b0aa314ce6ebf51b53313d4306cca8e46"
dependencies = [
"hermit-abi",
"libc",
"windows-sys 0.61.2",
"windows-sys 0.52.0",
]
[[package]]
@@ -1470,7 +1476,7 @@ dependencies = [
[[package]]
name = "metrics"
version = "3.3.0"
version = "4.0.2"
dependencies = [
"anyhow",
"log",
@@ -1498,7 +1504,7 @@ dependencies = [
[[package]]
name = "miner-cli"
version = "3.3.0"
version = "4.0.2"
dependencies = [
"clap",
"engine-cpu",
@@ -1509,13 +1515,14 @@ dependencies = [
"miner-service",
"num_cpus",
"primitive-types 0.13.1",
"quic-transport",
"rand 0.9.2",
"tokio",
]
[[package]]
name = "miner-service"
version = "3.3.0"
version = "4.0.2"
dependencies = [
"anyhow",
"crossbeam-channel",
@@ -1529,14 +1536,14 @@ dependencies = [
"pow-core",
"primitive-types 0.13.1",
"quantus-miner-api",
"quic-transport",
"quinn",
"rustls 0.21.12",
"tokio",
]
[[package]]
name = "miner-telemetry"
version = "3.3.0"
version = "4.0.2"
dependencies = [
"anyhow",
"futures",
@@ -1994,6 +2001,28 @@ dependencies = [
"plotters-backend",
]
[[package]]
name = "pool-service"
version = "4.0.2"
dependencies = [
"anyhow",
"clap",
"constant_time_eq",
"env_logger",
"hex",
"log",
"pow-core",
"primitive-types 0.13.1",
"quantus-miner-api",
"quic-transport",
"quinn",
"rand 0.9.2",
"serde",
"serde_json",
"tokio",
"warp",
]
[[package]]
name = "portable-atomic"
version = "1.13.1"
@@ -2020,10 +2049,11 @@ dependencies = [
[[package]]
name = "pow-core"
version = "3.3.0"
version = "4.0.2"
dependencies = [
"hex",
"primitive-types 0.13.1",
"qp-poseidon-core 3.1.0",
"qpow-math",
]
@@ -2231,6 +2261,12 @@ version = "2.1.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "78172860cd960773c72e97fba07ead9e5a1c13086c5746283744933e2e4a577b"
[[package]]
name = "qp-poseidon-core"
version = "3.1.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "5872607e25ea4ee5fb37e64bf1462168e1a36a4e719cdc8a105533c708253918"
[[package]]
name = "qpow-math"
version = "0.1.0"
@@ -2239,19 +2275,33 @@ dependencies = [
"hex",
"log",
"primitive-types 0.13.1",
"qp-poseidon-core",
"qp-poseidon-core 2.1.0",
]
[[package]]
name = "quantus-miner-api"
version = "0.2.0"
source = "git+https://github.com/Quantus-Network/chain.git?tag=v0.7.1-q-day-2#5ea94cb1e5347b12d091b6fa9f9234ceb897ec42"
version = "0.3.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "725bbf1de0bc7d2ace83f8e49ea754be5f45b84bc9471ddd3f417a9e2bb20553"
dependencies = [
"serde",
"serde_json",
"tokio",
]
[[package]]
name = "quic-transport"
version = "4.0.2"
dependencies = [
"anyhow",
"quantus-miner-api",
"quinn",
"rustls 0.21.12",
"serde_json",
"sha2",
"tokio",
]
[[package]]
name = "quinn"
version = "0.10.2"
@@ -2751,6 +2801,17 @@ dependencies = [
"digest",
]
[[package]]
name = "sha2"
version = "0.10.9"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "a7507d819769d01a365ab707794a4084392c824f54a7a6a7862f8c3d0892b283"
dependencies = [
"cfg-if",
"cpufeatures 0.2.17",
"digest",
]
[[package]]
name = "shlex"
version = "1.3.0"
@@ -2808,6 +2869,14 @@ dependencies = [
"windows-sys 0.61.2",
]
[[package]]
name = "solver-wasm"
version = "4.0.2"
dependencies = [
"primitive-types 0.13.1",
"qpow-math",
]
[[package]]
name = "spin"
version = "0.5.2"
@@ -3693,7 +3762,7 @@ version = "0.1.11"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "c2a7b1c03c876122aa43f3020e6c3c3ee5c05081c9a00739faf7503aeba10d22"
dependencies = [
"windows-sys 0.61.2",
"windows-sys 0.48.0",
]
[[package]]
Cargo.toml
+6 -3
View File
@@ -6,7 +6,10 @@ members = [
"crates/miner-cli",
"crates/miner-service",
"crates/miner-telemetry",
"crates/pool-service",
"crates/pow-core",
"crates/quic-transport",
"crates/solver-wasm",
]
default-members = ["crates/miner-cli"]
resolver = "2"
@@ -15,7 +18,7 @@ resolver = "2"
edition = "2021"
authors = ["Quantus Network"]
description = "Quantus External Miner Workspace"
version = "3.3.0"
version = "4.0.2"
[workspace.dependencies]
anyhow = "1"
@@ -29,9 +32,9 @@ num-bigint = { version = "0.4", features = ["rand"] }
num-traits = "0.2"
num_cpus = "1.16"
primitive-types = { version = "0.13.1", default-features = false }
qp-poseidon-core = { version = "2.1.0", default-features = false }
qp-poseidon-core = { version = "3.1.0", default-features = false }
qpow-math = { git = "https://github.com/Quantus-Network/chain.git", tag = "v0.7.1-q-day-2", package = "qpow-math", default-features = false }
quantus-miner-api = { git = "https://github.com/Quantus-Network/chain.git", tag = "v0.7.1-q-day-2" }
quantus-miner-api = "0.3.0"
rand = { version = "0.9", default-features = false, features = ["std", "std_rng", "thread_rng"] }
serde = { version = "1.0", features = ["derive"] }
serde_json = { version = "1.0.132", default-features = false }
Dockerfile
-54
View File
@@ -1,54 +0,0 @@
# syntax=docker/dockerfile:1
############################
# Builder stage
############################
FROM rust:1.85-slim-bookworm AS builder
# Install build dependencies
RUN apt-get update \
&& DEBIAN_FRONTEND=noninteractive apt-get install -y --no-install-recommends \
pkg-config \
libssl-dev \
ca-certificates \
&& rm -rf /var/lib/apt/lists/*
WORKDIR /build
# Copy workspace files
COPY Cargo.toml Cargo.lock rust-toolchain taplo.toml ./
COPY crates ./crates
# Build the miner-cli in release mode
RUN cargo build --release -p miner-cli --locked
# Strip debug symbols to reduce binary size
RUN strip target/release/quantus-miner || true
############################
# Runtime-only stage
############################
FROM debian:bookworm-slim
# Install runtime dependencies
RUN apt-get update \
&& DEBIAN_FRONTEND=noninteractive apt-get install -y --no-install-recommends \
ca-certificates \
&& rm -rf /var/lib/apt/lists/*
# Copy binary from builder stage
COPY --from=builder /build/target/release/quantus-miner /usr/local/bin/quantus-miner
# Expose miner API port and metrics port
EXPOSE 9833 9900
# Run as unprivileged user
RUN useradd --system --uid 10001 quantus
USER 10001:10001
# Default working directory
WORKDIR /data
# Start the miner
ENTRYPOINT ["quantus-miner"]
EXTERNAL_MINER_PROTOCOL.md
-235
View File
@@ -1,235 +0,0 @@
# External Miner Protocol Specification
This document defines the QUIC-based protocol for communication between the Quantus Network node and external QPoW miner services.
## Overview
The node delegates the mining task (finding a valid nonce) to external miner services over persistent QUIC connections. The node provides the necessary parameters (header hash, difficulty) and each external miner independently searches for a valid nonce according to the QPoW rules defined in the `qpow-math` crate. Miners push results back when found.
### Key Benefits of QUIC
- **Lower latency**: Results are pushed immediately when found (no polling)
- **Connection resilience**: Built-in connection migration and recovery
- **Multiplexed streams**: Multiple operations on single connection
- **Built-in TLS**: Encrypted by default
## Architecture
### Connection Model
```
┌─────────────────────────────────┐
│ Node │
│ (QUIC Server on port 9833) │
│ │
┌──────────┐ │ Broadcasts: NewJob │
│ Miner 1 │ ──connect───► │ Receives: JobResult │
└──────────┘ │ │
│ Supports multiple miners │
┌──────────┐ │ First valid result wins │
│ Miner 2 │ ──connect───► │ │
└──────────┘ └─────────────────────────────────┘
┌──────────┐
│ Miner 3 │ ──connect───►
└──────────┘
```
- **Node** acts as the QUIC server, listening on port 9833 (default)
- **Miners** act as QUIC clients, connecting to the node
- Single bidirectional stream per miner connection
- Connection persists across multiple mining jobs
- Multiple miners can connect simultaneously
### Multi-Miner Operation
When multiple miners are connected:
1. Node broadcasts the same `NewJob` to all connected miners
2. Each miner independently selects a random starting nonce
3. First miner to find a valid solution sends `JobResult`
4. Node uses the first valid result, ignores subsequent results for same job
5. New job broadcast implicitly cancels work on all miners
### Message Types
The protocol uses only **two message types**:
| Direction | Message | Description |
|-----------|---------|-------------|
| Node → Miner | `NewJob` | Submit a mining job (implicitly cancels any previous job) |
| Miner → Node | `JobResult` | Mining result (completed, failed, or cancelled) |
### Wire Format
Messages are length-prefixed JSON:
```
┌─────────────────┬─────────────────────────────────┐
│ Length (4 bytes)│ JSON payload (MinerMessage) │
│ big-endian u32 │ │
└─────────────────┴─────────────────────────────────┘
```
Maximum message size: 16 MB
## Data Types
See the `quantus-miner-api` crate for the canonical Rust definitions.
### MinerMessage (Enum)
```rust
pub enum MinerMessage {
NewJob(MiningRequest),
JobResult(MiningResult),
}
```
### MiningRequest
| Field | Type | Description |
|-------|------|-------------|
| `job_id` | String | Unique identifier (UUID recommended) |
| `mining_hash` | String | Header hash (64 hex chars, no 0x prefix) |
| `distance_threshold` | String | Difficulty (U512 as decimal string) |
Note: Nonce range is not specified - each miner independently selects a random starting point.
### MiningResult
| Field | Type | Description |
|-------|------|-------------|
| `status` | ApiResponseStatus | Result status (see below) |
| `job_id` | String | Job identifier |
| `nonce` | Option<String> | Winning nonce (U512 hex, no 0x prefix) |
| `work` | Option<String> | Winning nonce as bytes (128 hex chars) |
| `hash_count` | u64 | Number of nonces checked |
| `elapsed_time` | f64 | Time spent mining (seconds) |
### ApiResponseStatus (Enum)
| Value | Description |
|-------|-------------|
| `completed` | Valid nonce found |
| `failed` | Nonce range exhausted without finding solution |
| `cancelled` | Job was cancelled (new job received) |
| `running` | Job still in progress (not typically sent) |
## Protocol Flow
### Normal Mining Flow
```
Miner Node
│ │
│──── QUIC Connect ─────────────────────────►│
│◄─── Connection Established ────────────────│
│ │
│◄─── NewJob { job_id: "abc", ... } ─────────│
│ │
│ (picks random nonce, starts mining) │
│ │
│──── JobResult { job_id: "abc", ... } ─────►│ (found solution!)
│ │
│ (node submits block, gets new work) │
│ │
│◄─── NewJob { job_id: "def", ... } ─────────│
│ │
```
### Job Cancellation (Implicit)
When a new block arrives before the miner finds a solution, the node simply sends a new `NewJob`. The miner automatically cancels the previous job:
```
Miner Node
│ │
│◄─── NewJob { job_id: "abc", ... } ─────────│
│ │
│ (mining "abc") │
│ │
│ (new block arrives at node!) │
│ │
│◄─── NewJob { job_id: "def", ... } ─────────│
│ │
│ (cancels "abc", starts "def") │
│ │
│──── JobResult { job_id: "def", ... } ─────►│
```
### Miner Connect During Active Job
When a miner connects while a job is active, it immediately receives the current job:
```
Miner (new) Node
│ │ (already mining job "abc")
│──── QUIC Connect ─────────────────────────►│
│◄─── Connection Established ────────────────│
│◄─── NewJob { job_id: "abc", ... } ─────────│ (current job sent immediately)
│ │
│ (joins mining effort) │
```
### Stale Result Handling
If a result arrives for an old job, the node discards it:
```
Miner Node
│ │
│◄─── NewJob { job_id: "abc", ... } ─────────│
│ │
│◄─── NewJob { job_id: "def", ... } ─────────│ (almost simultaneous)
│ │
│──── JobResult { job_id: "abc", ... } ─────►│ (stale, node ignores)
│ │
│──── JobResult { job_id: "def", ... } ─────►│ (current, node uses)
```
## Configuration
### Node
```bash
# Listen for external miner connections on port 9833
quantus-node --miner-listen-port 9833
```
### Miner
```bash
# Connect to node
quantus-miner serve --node-addr 127.0.0.1:9833
```
## TLS Configuration
The node generates a self-signed TLS certificate at startup. The miner skips certificate verification by default (insecure mode). For production deployments, consider:
1. **Certificate pinning**: Configure the miner to accept only specific certificate fingerprints
2. **Proper CA**: Use certificates signed by a trusted CA
3. **Network isolation**: Run node and miner on a private network
## Error Handling
### Connection Loss
The miner automatically reconnects with exponential backoff:
- Initial delay: 1 second
- Maximum delay: 30 seconds
The node continues operating with remaining connected miners.
### Validation Errors
If the miner receives an invalid `MiningRequest`, it sends a `JobResult` with status `failed`.
## Notes
- All hex values should be sent **without** the `0x` prefix
- The miner implements validation logic from `qpow_math::is_valid_nonce`
- The node uses the `work` field from `MiningResult` to construct `QPoWSeal`
- ALPN protocol identifier: `quantus-miner`
- Each miner independently generates a random nonce starting point using cryptographically secure randomness
- With a 512-bit nonce space, collision between miners is statistically impossible
README.md
+59 -32
View File
@@ -16,18 +16,34 @@ The binary will be available at `target/release/quantus-miner`.
## Running
```bash
# CPU-only mining (default: auto-detected CPU cores)
./target/release/quantus-miner serve --cpu-workers 4
The node requires a shared auth token and TLS cert pin. Both live under the
node's chain config dir (`<base-path>/chains/<chain>/`):
# GPU-only mining
./target/release/quantus-miner serve --gpu-devices 1
- `miner-auth-token` — shared secret (**not** logged by the node; read this file)
- `miner-tls-cert-sha256` — SHA-256 of the miner TLS cert (also printed in node logs)
```bash
# Preferred: mount/read the node's chain config files
./target/release/quantus-miner serve \
--node-addr 127.0.0.1:9833 \
--auth-token-file /path/to/miner-auth-token \
--tls-cert-sha256-file /path/to/miner-tls-cert-sha256 \
--cpu-workers 4
# Or pass values directly (token from miner-auth-token; fingerprint also in node logs)
./target/release/quantus-miner serve \
--node-addr 127.0.0.1:9833 \
--auth-token <TOKEN> \
--tls-cert-sha256 <FINGERPRINT> \
--gpu-devices 1
# Hybrid CPU+GPU mining
./target/release/quantus-miner serve --cpu-workers 4 --gpu-devices 1
# Custom port and metrics
./target/release/quantus-miner serve --cpu-workers 2 --port 8000 --metrics-port 9900
./target/release/quantus-miner serve \
--node-addr 127.0.0.1:9833 \
--auth-token-file /path/to/miner-auth-token \
--tls-cert-sha256-file /path/to/miner-tls-cert-sha256 \
--cpu-workers 4 \
--gpu-devices 1
```
## Configuration
@@ -35,6 +51,10 @@ The binary will be available at `target/release/quantus-miner`.
| Argument | Environment Variable | Description | Default |
|----------|---------------------|-------------|---------|
| `--node-addr <ADDR>` | `MINER_NODE_ADDR` | Node address to connect to | `127.0.0.1:9833` |
| `--auth-token <TOKEN>` | `MINER_AUTH_TOKEN` | Shared secret from the node's `miner-auth-token` file (not logged) | required |
| `--auth-token-file <PATH>` | `MINER_AUTH_TOKEN_FILE` | Read the shared secret from a file (preferred) | — |
| `--tls-cert-sha256 <HEX>` | `MINER_TLS_CERT_SHA256` | SHA-256 of the node's miner TLS cert (`miner-tls-cert-sha256` / node logs) | required |
| `--tls-cert-sha256-file <PATH>` | `MINER_TLS_CERT_SHA256_FILE` | Read the TLS cert fingerprint from a file | — |
| `--cpu-workers <N>` | `MINER_CPU_WORKERS` | Number of CPU worker threads | Auto-detect |
| `--gpu-devices <N>` | `MINER_GPU_DEVICES` | Number of GPU devices | Auto-detect |
| `--gpu-batch-size <N>` | `MINER_GPU_BATCH_SIZE` | GPU batch size in nonces | 1000000 |
@@ -70,27 +90,46 @@ cargo build -p miner-cli --release
## Examples
All `serve` examples need the auth token and TLS pin (files or inline values).
```bash
# CPU mining with 8 workers
./target/release/quantus-miner serve --cpu-workers 8
./target/release/quantus-miner serve \
--auth-token-file /path/to/miner-auth-token \
--tls-cert-sha256-file /path/to/miner-tls-cert-sha256 \
--cpu-workers 8
# Pure GPU mining
./target/release/quantus-miner serve --gpu-devices 1
./target/release/quantus-miner serve \
--auth-token-file /path/to/miner-auth-token \
--tls-cert-sha256-file /path/to/miner-tls-cert-sha256 \
--gpu-devices 1
# GPU mining with throttle (reduce GPU utilization)
./target/release/quantus-miner serve --gpu-devices 1 --gpu-throttle-ms 50
./target/release/quantus-miner serve \
--auth-token-file /path/to/miner-auth-token \
--tls-cert-sha256-file /path/to/miner-tls-cert-sha256 \
--gpu-devices 1 --gpu-throttle-ms 50
# Hybrid mining: 4 CPU + 1 GPU workers
./target/release/quantus-miner serve --cpu-workers 4 --gpu-devices 1
./target/release/quantus-miner serve \
--auth-token-file /path/to/miner-auth-token \
--tls-cert-sha256-file /path/to/miner-tls-cert-sha256 \
--cpu-workers 4 --gpu-devices 1
# With verbose logging
RUST_LOG=debug ./target/release/quantus-miner serve --cpu-workers 2 --gpu-devices 1
RUST_LOG=debug ./target/release/quantus-miner serve \
--auth-token-file /path/to/miner-auth-token \
--tls-cert-sha256-file /path/to/miner-tls-cert-sha256 \
--cpu-workers 2 --gpu-devices 1
# Production setup with metrics
./target/release/quantus-miner serve \
--node-addr 127.0.0.1:9833 \
--auth-token-file /path/to/miner-auth-token \
--tls-cert-sha256-file /path/to/miner-tls-cert-sha256 \
--cpu-workers 6 \
--gpu-devices 1 \
--port 9833 \
--metrics-port 9900
```
@@ -98,25 +137,13 @@ RUST_LOG=debug ./target/release/quantus-miner serve --cpu-workers 2 --gpu-device
The miner uses a QUIC-based protocol for communication with the node:
- **Transport**: QUIC with TLS 1.3 (self-signed certificates)
- **Transport**: QUIC with TLS 1.3 (self-signed certificate, pinned by SHA-256)
- **Auth**: `Ready { token }` must match the node's `miner-auth-token`
- **ALPN**: `quantus-miner/2`
- **Port**: 9833 (default)
- **Messages**: `NewJob` (from node) and `JobResult` (from miner)
- **Messages**: `Ready` (miner→node), `NewJob` (node→miner), `JobResult` (miner→node)
For full protocol specification, see `EXTERNAL_MINER_PROTOCOL.md`.
## Docker
```bash
# Quick start
docker pull ghcr.io/quantus-network/quantus-miner:latest
docker run -d -p 9833:9833 -p 9900:9900 \
ghcr.io/quantus-network/quantus-miner:latest \
--cpu-workers 4 --metrics-port 9900
# Build from source
docker build -t quantus-miner .
docker run -d -p 9833:9833 quantus-miner serve --cpu-workers 4
```
For full protocol specification, see the node's `MINING.md`.
## Benchmarking
SECURITY.md
+80
View File
@@ -0,0 +1,80 @@
# Security Policy
Quantus Network takes the security of our users seriously and welcomes reports
from security researchers.
## Reporting a vulnerability
**Please do not open a public issue, pull request, or social-media post for a
security vulnerability.** Public disclosure before a fix is available puts
users' funds at risk. Use one of the private channels below and we will
coordinate a fix and disclosure with you.
1. **GitHub private security advisory (preferred).**
[Report a vulnerability](https://github.com/quantus-network/quantus-miner/security/advisories/new).
This gives you a private, structured thread with the maintainers and is the
fastest way to reach us.
2. **Email.** Send details to **security@quantus.com**.
### What to include
- A clear description of the issue and its security impact.
- A working proof of concept is required. It must run directly against this
project, using a release build or locally built version, and demonstrate the
reported behavior and security impact. Standalone code, mathematical
examples, or simulations that only reproduce the theory without exercising
the project do not satisfy this requirement.
- Step-by-step reproduction instructions.
- Affected platforms and versions.
- Any relevant logs, addresses, or transaction IDs (for on-chain issues).
If you used AI tooling to find or write up the report, please say so.
## Our commitment (safe harbor)
We consider security research conducted in good faith under this policy to be
authorized. We will not pursue or support legal action against researchers who:
- make a good-faith effort to avoid privacy violations, data destruction, and
interruption or degradation of our services;
- only interact with accounts they own or have explicit permission to access;
and
- give us a reasonable opportunity to fix an issue before disclosing it
publicly.
If in doubt about whether an action is authorized, ask us first at
security@quantus.com.
## What to expect
- **Acknowledgement:** within **2 business days**.
- **Triage and initial assessment:** within **7 business days**.
- **Coordinated disclosure:** we aim to ship a fix and coordinate public
disclosure within **90 days** of the report. We will keep you updated on
progress and agree on a disclosure date with you.
- **Credit:** with your permission, we are happy to publicly credit you for the
report once a fix is released.
## Scope
**In scope:** the external mining service (CPU and GPU) in this repository — anything that could lead to theft or misattribution of mining rewards, invalid blocks being produced, or remote code execution or crash via its service API.
**Out of scope:** issues in third-party services or nodes we do not operate;
reports generated solely by automated scanners without a demonstrated impact;
low-severity or informational issues on our marketing and landing websites;
and social-engineering or physical attacks.
## Rewards
At our **sole discretion**, we may offer a reward for a valid report. To be
eligible, a report must be submitted **privately** through one of the channels
above and identify a genuine vulnerability with real impact on users. Trivial
or low-impact findings, automated-scanner output without a working proof of
concept, and already-known issues are **not** eligible. There is no fixed
bounty and no guaranteed payout; whether a report qualifies, and any amount,
are determined solely by Quantus Network.
## Supported versions
Only the latest release is supported; security fixes are delivered in new
versions. Please keep your installation up to date.
clippy.sh
Executable
+7
View File
@@ -0,0 +1,7 @@
#!/usr/bin/env bash
# Run the same clippy command as CI (see .github/workflows/ci.yml)
set -euo pipefail
cargo +nightly fmt
taplo format
cargo clippy --workspace --all-targets --all-features --locked -- -D warnings
crates/engine-cpu/src/lib.rs
+4
View File
@@ -51,6 +51,10 @@ pub enum EngineStatus {
Cancelled {
hash_count: u64,
},
/// GPU device lost or unresponsive. Worker should exit permanently.
DeviceLost {
hash_count: u64,
},
}
/// Cancellation checker passed to search_range.
crates/engine-gpu/benches/gpu_engine_bench.rs
+20 -6
View File
@@ -6,9 +6,17 @@ use primitive_types::U512;
use rand::RngCore;
use std::sync::atomic::AtomicBool;
/// Drop thread-local wgpu buffers before `GpuEngine` is dropped. Criterion
/// creates a fresh engine per group; without this, TLS buffers outlive the
/// device and the next group panics (`Buffer[…] does not exist`).
fn teardown_gpu(engine: GpuEngine) {
GpuEngine::clear_worker_resources();
drop(engine);
}
fn bench_cpu_vs_gpu_small(c: &mut Criterion) {
let cpu_engine = FastCpuEngine::new(10_000);
let gpu_engine = GpuEngine::try_new(10_000_000, 0).expect("Failed to init GPU");
let gpu_engine = GpuEngine::try_new(10_000_000, 0, false).expect("Failed to init GPU");
let cancel_flag = AtomicBool::new(false);
let cancel_check = AtomicBoolCancelCheck(&cancel_flag);
@@ -55,11 +63,12 @@ fn bench_cpu_vs_gpu_small(c: &mut Criterion) {
});
group.finish();
teardown_gpu(gpu_engine);
}
fn bench_cpu_vs_gpu_medium(c: &mut Criterion) {
let cpu_engine = FastCpuEngine::new(10_000);
let gpu_engine = GpuEngine::try_new(10_000_000, 0).expect("Failed to init GPU");
let gpu_engine = GpuEngine::try_new(10_000_000, 0, false).expect("Failed to init GPU");
let cancel_flag = AtomicBool::new(false);
let cancel_check = AtomicBoolCancelCheck(&cancel_flag);
@@ -106,11 +115,12 @@ fn bench_cpu_vs_gpu_medium(c: &mut Criterion) {
});
group.finish();
teardown_gpu(gpu_engine);
}
fn bench_cpu_vs_gpu_large(c: &mut Criterion) {
let cpu_engine = FastCpuEngine::new(10_000);
let gpu_engine = GpuEngine::try_new(10_000_000, 0).expect("Failed to init GPU");
let gpu_engine = GpuEngine::try_new(10_000_000, 0, false).expect("Failed to init GPU");
let cancel_flag = AtomicBool::new(false);
let cancel_check = AtomicBoolCancelCheck(&cancel_flag);
@@ -157,11 +167,12 @@ fn bench_cpu_vs_gpu_large(c: &mut Criterion) {
});
group.finish();
teardown_gpu(gpu_engine);
}
fn bench_solution_finding(c: &mut Criterion) {
let cpu_engine = FastCpuEngine::new(10_000);
let gpu_engine = GpuEngine::try_new(10_000_000, 0).expect("Failed to init GPU");
let gpu_engine = GpuEngine::try_new(10_000_000, 0, false).expect("Failed to init GPU");
let cancel_flag = AtomicBool::new(false);
let cancel_check = AtomicBoolCancelCheck(&cancel_flag);
@@ -208,11 +219,12 @@ fn bench_solution_finding(c: &mut Criterion) {
});
group.finish();
teardown_gpu(gpu_engine);
}
fn bench_throughput_per_second(c: &mut Criterion) {
let cpu_engine = FastCpuEngine::new(10_000);
let gpu_engine = GpuEngine::try_new(10_000_000, 0).expect("Failed to init GPU");
let gpu_engine = GpuEngine::try_new(10_000_000, 0, false).expect("Failed to init GPU");
let cancel_flag = AtomicBool::new(false);
let cancel_check = AtomicBoolCancelCheck(&cancel_flag);
@@ -259,10 +271,11 @@ fn bench_throughput_per_second(c: &mut Criterion) {
});
group.finish();
teardown_gpu(gpu_engine);
}
fn bench_gpu_batch_efficiency(c: &mut Criterion) {
let gpu_engine = GpuEngine::try_new(10_000_000, 0).expect("Failed to init GPU");
let gpu_engine = GpuEngine::try_new(10_000_000, 0, false).expect("Failed to init GPU");
let cancel_flag = AtomicBool::new(false);
let cancel_check = AtomicBoolCancelCheck(&cancel_flag);
@@ -335,6 +348,7 @@ fn bench_gpu_batch_efficiency(c: &mut Criterion) {
});
group.finish();
teardown_gpu(gpu_engine);
}
criterion_group!(
crates/engine-gpu/examples/gpu_cpu_parity.rs
+63
View File
@@ -0,0 +1,63 @@
use engine_cpu::{AtomicBoolCancelCheck, EngineStatus, MinerEngine, Range};
use engine_gpu::GpuEngine;
use primitive_types::U512;
use rand::RngCore;
use std::sync::atomic::AtomicBool;
fn main() {
env_logger::init();
let jobs: usize = std::env::args()
.nth(1)
.map(|s| s.parse().expect("job count"))
.unwrap_or(25);
let engine = GpuEngine::try_new(1_000_000, 0, false).expect("GPU init failed");
let cancel_flag = AtomicBool::new(false);
let cancel = AtomicBoolCancelCheck(&cancel_flag);
let mut rng = rand::rng();
let mut found = 0usize;
for job in 0..jobs {
let mut header = [0u8; 32];
rng.fill_bytes(&mut header);
let difficulty = U512::from(100_000u64);
let ctx = engine.prepare_context(header, difficulty);
let start = if job == 0 {
// Cross a 2^256 boundary: the high nonce half changes mid-range,
// exercising the midstate batch clamp.
(U512::from(3u64) << 256) - U512::from(1_000u64)
} else {
let mut start_bytes = [0u8; 64];
rng.fill_bytes(&mut start_bytes);
// Keep clear of the very top so range arithmetic cannot wrap
start_bytes[0] = 0;
U512::from_big_endian(&start_bytes)
};
let range = Range {
start,
end: start + U512::from(10_000_000u64),
};
match engine.search_range(&ctx, range.clone(), &cancel) {
EngineStatus::Found { candidate, .. } => {
let cpu_hash = pow_core::hash_from_nonce(&ctx, candidate.nonce);
assert_eq!(
cpu_hash, candidate.hash,
"job {job}: GPU hash != CPU hash for nonce {}",
candidate.nonce
);
assert!(cpu_hash < ctx.target, "job {job}: hash not below target");
assert!(
candidate.nonce >= range.start && candidate.nonce <= range.end,
"job {job}: nonce outside range"
);
found += 1;
}
EngineStatus::Exhausted { .. } => {}
other => panic!("job {job}: unexpected status {other:?}"),
}
}
assert!(found > 0, "no solutions found across {jobs} jobs");
println!("PARITY OK: {found}/{jobs} jobs found solutions, all verified against CPU");
}
crates/engine-gpu/examples/gpu_features.rs
+30
View File
@@ -0,0 +1,30 @@
fn main() {
let instance = wgpu::Instance::new(&wgpu::InstanceDescriptor {
backends: wgpu::Backends::PRIMARY,
..Default::default()
});
for adapter in instance.enumerate_adapters(wgpu::Backends::PRIMARY) {
let info = adapter.get_info();
let features = adapter.features();
println!("{} ({:?}, {:?})", info.name, info.device_type, info.backend);
println!(
" SHADER_INT64: {}",
features.contains(wgpu::Features::SHADER_INT64)
);
println!(
" SUBGROUP: {}",
features.contains(wgpu::Features::SUBGROUP)
);
println!(
" TIMESTAMP_QUERY: {}",
features.contains(wgpu::Features::TIMESTAMP_QUERY)
);
let limits = adapter.limits();
println!(
" max_workgroup_size_x: {}, max_invocations: {}, max_workgroups_per_dim: {}",
limits.max_compute_workgroup_size_x,
limits.max_compute_invocations_per_workgroup,
limits.max_compute_workgroups_per_dimension
);
}
}
crates/engine-gpu/examples/hashrate.rs
+52
View File
@@ -0,0 +1,52 @@
use engine_cpu::{AtomicBoolCancelCheck, MinerEngine, Range};
use engine_gpu::GpuEngine;
use primitive_types::U512;
use std::sync::atomic::AtomicBool;
fn main() {
env_logger::init();
let args: Vec<String> = std::env::args().collect();
let total: u64 = args
.get(1)
.map(|s| s.parse().expect("total nonces"))
.unwrap_or(4_000_000);
let batch: u32 = args
.get(2)
.map(|s| s.parse().expect("batch size"))
.unwrap_or(1_000_000);
let engine = GpuEngine::try_new(batch, 0, false).expect("GPU init failed");
let cancel_flag = AtomicBool::new(false);
let cancel = AtomicBoolCancelCheck(&cancel_flag);
let header = [42u8; 32];
let ctx = engine.prepare_context(header, U512::from(u64::MAX));
let range = Range {
start: U512::from(1u64) << 200,
end: (U512::from(1u64) << 200) + U512::from(total - 1),
};
// Warmup
let warm = Range {
start: range.start,
end: range.start + U512::from(200_000u64),
};
engine.search_range(&ctx, warm, &cancel);
let start = std::time::Instant::now();
let status = engine.search_range(&ctx, range, &cancel);
let elapsed = start.elapsed().as_secs_f64();
let hashes = match status {
engine_cpu::EngineStatus::Exhausted { hash_count } => hash_count,
engine_cpu::EngineStatus::Found { hash_count, .. } => hash_count,
other => panic!("unexpected status: {other:?}"),
};
println!(
"{} hashes in {:.3}s = {:.3} MH/s",
hashes,
elapsed,
hashes as f64 / elapsed / 1e6
);
}
crates/engine-gpu/examples/trusted_hashrate.rs
+229
View File
@@ -0,0 +1,229 @@
use pow_core::{hash_from_nonce, mining_midstate, JobContext};
use primitive_types::U512;
use std::time::Instant;
fn u512_to_u32s_le(v: U512) -> [u32; 16] {
let bytes = v.to_little_endian();
let mut out = [0u32; 16];
for i in 0..16 {
out[i] = u32::from_le_bytes(bytes[i * 4..(i + 1) * 4].try_into().unwrap());
}
out
}
struct Runner {
device: wgpu::Device,
queue: wgpu::Queue,
pipeline: wgpu::ComputePipeline,
results: wgpu::Buffer,
midstate: wgpu::Buffer,
start_nonce: wgpu::Buffer,
target: wgpu::Buffer,
cfg: wgpu::Buffer,
staging: wgpu::Buffer,
}
impl Runner {
async fn new(trusted: bool) -> Self {
let instance = wgpu::Instance::new(&wgpu::InstanceDescriptor {
backends: wgpu::Backends::PRIMARY,
..Default::default()
});
let adapter = instance
.request_adapter(&wgpu::RequestAdapterOptions::default())
.await
.expect("no adapter");
assert!(
adapter.features().contains(wgpu::Features::SHADER_INT64),
"SHADER_INT64 required"
);
let (device, queue) = adapter
.request_device(&wgpu::DeviceDescriptor {
label: None,
required_features: wgpu::Features::SHADER_INT64,
..Default::default()
})
.await
.unwrap();
let kernel = engine_gpu::Kernel::for_adapter(&adapter);
let desc = wgpu::ShaderModuleDescriptor {
label: Some(kernel.label()),
source: wgpu::ShaderSource::Wgsl(kernel.source().into()),
};
let shader = if trusted {
unsafe {
device.create_shader_module_trusted(desc, wgpu::ShaderRuntimeChecks::unchecked())
}
} else {
device.create_shader_module(desc)
};
let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
label: None,
layout: None,
module: &shader,
entry_point: Some("mining_main"),
compilation_options: Default::default(),
cache: None,
});
let mk = |size: u64, usage| {
device.create_buffer(&wgpu::BufferDescriptor {
label: None,
size,
usage,
mapped_at_creation: false,
})
};
use wgpu::BufferUsages as U;
Runner {
pipeline,
results: mk(132, U::STORAGE | U::COPY_SRC | U::COPY_DST),
midstate: mk(96, U::STORAGE | U::COPY_DST),
start_nonce: mk(64, U::STORAGE | U::COPY_DST),
target: mk(64, U::STORAGE | U::COPY_DST),
cfg: mk(12, U::STORAGE | U::COPY_DST),
staging: mk(132, U::MAP_READ | U::COPY_DST),
device,
queue,
}
}
fn run_batch(&self, header: [u8; 32], start: U512, batch: u32, target: U512) -> f64 {
let nonce_be = start.to_big_endian();
let mid = mining_midstate(header, nonce_be[..32].try_into().unwrap());
let mut mid_u32 = [0u32; 24];
for (i, f) in mid.iter().enumerate() {
mid_u32[2 * i] = *f as u32;
mid_u32[2 * i + 1] = (*f >> 32) as u32;
}
self.queue
.write_buffer(&self.midstate, 0, bytemuck::cast_slice(&mid_u32));
self.queue.write_buffer(
&self.start_nonce,
0,
bytemuck::cast_slice(&u512_to_u32s_le(start)),
);
self.queue.write_buffer(
&self.target,
0,
bytemuck::cast_slice(&u512_to_u32s_le(target)),
);
self.queue
.write_buffer(&self.cfg, 0, bytemuck::cast_slice(&[batch, 1u32, batch]));
self.queue.write_buffer(&self.results, 0, &[0u8; 132]);
let layout = self.pipeline.get_bind_group_layout(0);
let entries: Vec<wgpu::BindGroupEntry<'_>> = [
(0, &self.results),
(1, &self.midstate),
(2, &self.start_nonce),
(3, &self.target),
(4, &self.cfg),
]
.iter()
.map(|(i, b)| wgpu::BindGroupEntry {
binding: *i,
resource: b.as_entire_binding(),
})
.collect();
let bind_group = self.device.create_bind_group(&wgpu::BindGroupDescriptor {
label: None,
layout: &layout,
entries: &entries,
});
let t0 = Instant::now();
let mut encoder = self
.device
.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
{
let mut cpass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
label: None,
timestamp_writes: None,
});
cpass.set_pipeline(&self.pipeline);
cpass.set_bind_group(0, &bind_group, &[]);
cpass.dispatch_workgroups(batch.div_ceil(256), 1, 1);
}
encoder.copy_buffer_to_buffer(&self.results, 0, &self.staging, 0, 132);
self.queue.submit(Some(encoder.finish()));
let slice = self.staging.slice(..);
let (tx, rx) = std::sync::mpsc::channel();
slice.map_async(wgpu::MapMode::Read, move |r| tx.send(r).unwrap());
loop {
let _ = self.device.poll(wgpu::PollType::Wait {
submission_index: None,
timeout: None,
});
if rx.try_recv().is_ok() {
break;
}
}
let elapsed = t0.elapsed().as_secs_f64();
self.staging.unmap();
elapsed
}
fn read_results(&self) -> [u32; 33] {
let slice = self.staging.slice(..);
let (tx, rx) = std::sync::mpsc::channel();
slice.map_async(wgpu::MapMode::Read, move |r| tx.send(r).unwrap());
let _ = self.device.poll(wgpu::PollType::Wait {
submission_index: None,
timeout: None,
});
rx.recv().unwrap().unwrap();
let data = slice.get_mapped_range();
let mut out = [0u32; 33];
out.copy_from_slice(bytemuck::cast_slice(&data));
drop(data);
self.staging.unmap();
out
}
}
fn main() {
let batch: u32 = std::env::args()
.nth(1)
.map(|s| s.parse().unwrap())
.unwrap_or(8_000_000);
let iters = 4u32;
let rt = tokio::runtime::Runtime::new().unwrap();
for trusted in [false, true] {
let label = if trusted {
"trusted (no checks)"
} else {
"checked (default)"
};
let runner = rt.block_on(Runner::new(trusted));
let header = [9u8; 32];
let ctx = JobContext::new(header, U512::one());
let start = (U512::from(7u64) << 300) | U512::from(123456789u64);
runner.run_batch(header, start, 256, ctx.target);
let r = runner.read_results();
assert_eq!(r[0], 1, "{label}: no solution with target=MAX");
let nonce = U512::from_little_endian(bytemuck::cast_slice(&r[1..17]));
let hash = U512::from_little_endian(bytemuck::cast_slice(&r[17..33]));
assert_eq!(hash, hash_from_nonce(&ctx, nonce), "{label}: hash mismatch");
println!("{label}: correctness OK");
let header = [42u8; 32];
let start = U512::from(1u64) << 200;
let target = U512::one();
runner.run_batch(header, start, batch, target);
let mut total = 0.0;
for i in 0..iters {
total += runner.run_batch(
header,
start + U512::from((i as u64 + 1) * batch as u64),
batch,
target,
);
}
let mhs = (batch as f64 * iters as f64) / total / 1e6;
println!("{label}: {batch} nonces x{iters} in {total:.3}s = {mhs:.3} MH/s");
}
}
crates/engine-gpu/examples/verify_nonce.rs
+4 -1
View File
@@ -26,7 +26,7 @@ fn main() {
// 3. Verify with GPU engine
log::info!("Initializing GPU engine...");
let gpu_engine = GpuEngine::try_new(10_000_000, 0).expect("Failed to init GPU");
let gpu_engine = GpuEngine::try_new(10_000_000, 0, false).expect("Failed to init GPU");
// Search a small range around the valid nonce
let gpu_range = Range {
@@ -56,6 +56,9 @@ fn main() {
EngineStatus::Cancelled { .. } => {
log::error!("FAILURE: GPU search cancelled!");
}
EngineStatus::DeviceLost { .. } => {
log::error!("FAILURE: GPU device lost!");
}
EngineStatus::Running { .. } => {
log::error!("FAILURE: GPU returned Running status!");
}
crates/engine-gpu/src/end_to_end_tests.rs
+16 -33
View File
@@ -7,6 +7,7 @@ use wgpu::util::DeviceExt;
pub async fn test_end_to_end_mining(
device: &wgpu::Device,
queue: &wgpu::Queue,
shader_src: &str,
) -> Result<(), Box<dyn std::error::Error>> {
println!("Running End-to-End Mining Test...");
@@ -28,15 +29,17 @@ pub async fn test_end_to_end_mining(
// 4. Run GPU Mining for this specific nonce
// Header Buffer
let mut header_u32s = [0u32; 8];
for (i, item) in header_u32s.iter_mut().enumerate() {
let chunk = &ctx.header[i * 4..(i + 1) * 4];
*item = u32::from_le_bytes([chunk[0], chunk[1], chunk[2], chunk[3]]);
// Midstate Buffer (header + high nonce half absorbed on CPU)
let nonce_be = nonce_val.to_big_endian();
let midstate = pow_core::mining_midstate(ctx.header, nonce_be[..32].try_into().unwrap());
let mut midstate_u32s = [0u32; 24];
for (i, felt) in midstate.iter().enumerate() {
midstate_u32s[2 * i] = *felt as u32;
midstate_u32s[2 * i + 1] = (*felt >> 32) as u32;
}
let header_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
label: Some("Header Buffer"),
contents: bytemuck::cast_slice(&header_u32s),
let midstate_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
label: Some("Midstate Buffer"),
contents: bytemuck::cast_slice(&midstate_u32s),
usage: wgpu::BufferUsages::STORAGE,
});
@@ -82,11 +85,11 @@ pub async fn test_end_to_end_mining(
let zeros = vec![0u8; results_size];
queue.write_buffer(&results_buffer, 0, &zeros);
// Dispatch config buffer: [total_threads, nonces_per_thread, total_nonces, cancel_check_interval]
let dispatch_config_data: [u32; 4] = [256, 1, 256, 10000];
// Dispatch config buffer: [total_threads, nonces_per_thread, total_nonces]
let dispatch_config_data: [u32; 3] = [256, 1, 256];
let dispatch_config_buffer = device.create_buffer(&wgpu::BufferDescriptor {
label: Some("Dispatch Config Buffer"),
size: 16,
size: 12,
usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
mapped_at_creation: false,
});
@@ -96,25 +99,9 @@ pub async fn test_end_to_end_mining(
bytemuck::cast_slice(&dispatch_config_data),
);
// Cancel flag buffer: 0 = running, 1 = cancel requested
let cancel_flag_data: [u32; 1] = [0];
let cancel_flag_buffer = device.create_buffer(&wgpu::BufferDescriptor {
label: Some("Cancel Flag Buffer"),
size: 4,
usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
mapped_at_creation: false,
});
queue.write_buffer(
&cancel_flag_buffer,
0,
bytemuck::cast_slice(&cancel_flag_data),
);
// Load Shader
let shader_source = include_str!("mining.wgsl");
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: Some("Mining Shader"),
source: wgpu::ShaderSource::Wgsl(shader_source.into()),
source: wgpu::ShaderSource::Wgsl(shader_src.into()),
});
let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
@@ -137,7 +124,7 @@ pub async fn test_end_to_end_mining(
},
wgpu::BindGroupEntry {
binding: 1,
resource: header_buffer.as_entire_binding(),
resource: midstate_buffer.as_entire_binding(),
},
wgpu::BindGroupEntry {
binding: 2,
@@ -151,10 +138,6 @@ pub async fn test_end_to_end_mining(
binding: 4,
resource: dispatch_config_buffer.as_entire_binding(),
},
wgpu::BindGroupEntry {
binding: 5,
resource: cancel_flag_buffer.as_entire_binding(),
},
],
});
crates/engine-gpu/src/gpu_tiers.rs
+25
View File
@@ -423,6 +423,31 @@ const QUALCOMM_TIERS: &[GpuTier] = &[
];
const APPLE_TIERS: &[GpuTier] = &[
// M5 series
GpuTier {
pattern: r"m5 ultra",
name: "Apple M5 Ultra",
workgroup_divisor: 4,
min_workgroups: 1600,
},
GpuTier {
pattern: r"m5 max",
name: "Apple M5 Max",
workgroup_divisor: 4,
min_workgroups: 800,
},
GpuTier {
pattern: r"m5 pro",
name: "Apple M5 Pro",
workgroup_divisor: 4,
min_workgroups: 400,
},
GpuTier {
pattern: r"\bm5\b",
name: "Apple M5",
workgroup_divisor: 4,
min_workgroups: 200,
},
// M4 series
GpuTier {
pattern: r"m4 ultra",
crates/engine-gpu/src/mining.wgsl → crates/engine-gpu/src/kernels/mining.wgsl
+64 -42
View File
@@ -191,7 +191,9 @@ const MDS_MATRIX_DIAG_12: array<array<u32, 2>, 12> = array<array<u32, 2>, 12>(
// Storage buffers
@group(0) @binding(0) var<storage, read_write> results: array<atomic<u32>>;
@group(0) @binding(1) var<storage, read> header: array<u32, 8>; // 32 bytes
// 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>; // 64 bytes
@group(0) @binding(3) var<storage, read> difficulty_target: array<u32, 16>; // 64 bytes (U512 target)
@group(0) @binding(4) var<storage, read> dispatch_config: array<u32, 3>; // [total_threads, nonces_per_thread, total_nonces]
@@ -787,17 +789,8 @@ fn hash_squeeze_twice(input: array<u32, 24>) -> array<u32, 16> {
return poseidon2_hash_squeeze_twice(input);
}
// Check if hash < target (U512 comparison)
fn is_below_target(hash: array<u32, 16>, difficulty_tgt: array<u32, 16>) -> bool {
// Compare from most significant to least significant
for (var i = 0u; i < 16u; i++) {
if (hash[15u - i] < difficulty_tgt[15u - i]) {
return true;
} else if (hash[15u - i] > difficulty_tgt[15u - i]) {
return false;
}
}
return false; // Equal, not below
fn bswap32(v: u32) -> u32 {
return ((v & 0xFFu) << 24u) | ((v & 0xFF00u) << 8u) | ((v >> 8u) & 0xFF00u) | (v >> 24u);
}
// Main mining kernel
@@ -834,61 +827,90 @@ fn mining_main(@builtin(global_invocation_id) global_id: vec3<u32>) {
return;
}
// current_nonce = start_nonce + logical_index
// current_nonce = start_nonce + logical_index. The host guarantees a batch
// never carries into the high nonce half (limbs 8..15).
var current_nonce: array<u32, 16>;
var carry: u32 = 0u;
// Add logical_index into the low limb and propagate carry through the U512 nonce
let val0 = start_nonce[0];
let sum0 = val0 + logical_index;
current_nonce[0] = sum0;
carry = select(0u, 1u, sum0 < val0);
// Propagate carry through remaining limbs
for (var i = 1u; i < 16u; i++) {
for (var i = 1u; i < 8u; i++) {
let val = start_nonce[i];
let sum = val + carry;
current_nonce[i] = sum;
carry = select(0u, 1u, sum < val);
}
// Construct input (96 bytes = 24 u32s)
// Header (32 bytes = 8 u32s) followed by Nonce (64 bytes = 16 u32s)
var input: array<u32, 24>;
for (var i = 0u; i < 8u; i++) {
input[i] = header[i];
for (var i = 8u; i < 16u; i++) {
current_nonce[i] = start_nonce[i];
}
// Nonce needs to be Big Endian in the byte stream for hashing.
// current_nonce is Little Endian words.
for (var i = 0u; i < 16u; i++) {
let val = current_nonce[15u - i];
// Reverse bytes
// Resume the sponge from the precomputed midstate: absorb the low nonce
// half (as Big Endian words), pad, squeeze twice (3 permutations instead of 5).
var state: array<GoldilocksField, 12>;
for (var i = 0u; i < 12u; i++) {
state[i] = gf_from_limbs(midstate[2u * i], midstate[2u * i + 1u]);
}
for (var i = 0u; i < 8u; i++) {
let val = current_nonce[7u - i];
let rev = ((val & 0xFFu) << 24u) |
((val & 0xFF00u) << 8u) |
((val & 0xFF0000u) >> 8u) |
((val & 0xFF000000u) >> 24u);
input[8u + i] = rev;
state[i] = gf_add(state[i], gf_from_u32(rev));
}
poseidon2_permute(&state);
state[0] = gf_add(state[0], gf_one());
state[1] = gf_add(state[1], gf_one());
poseidon2_permute(&state);
// Hash (Big Endian)
let hash_be = hash_squeeze_twice(input);
// Convert to Little Endian for difficulty check and storage
// 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.
let first_output = field_elements_to_bytes(array<GoldilocksField, 4>(
state[0], state[1], state[2], state[3]
));
var hash_le: array<u32, 16>;
for (var i = 0u; i < 16u; i++) {
let val = hash_be[15u - i];
// Reverse bytes in u32
hash_le[i] = ((val & 0xFFu) << 24u) |
((val & 0xFF00u) << 8u) |
((val & 0xFF0000u) >> 8u) |
((val & 0xFF000000u) >> 24u);
for (var i = 0u; i < 8u; i++) {
hash_le[15u - i] = bswap32(first_output[i]);
}
var cmp = 0u;
for (var i = 0u; i < 8u; i++) {
let h = hash_le[15u - i];
let t = difficulty_target[15u - i];
if (h != t) {
cmp = select(2u, 1u, h > t);
break;
}
}
if (cmp == 1u) {
continue;
}
// Check target
if (is_below_target(hash_le, difficulty_target)) {
poseidon2_permute(&state);
let second_output = field_elements_to_bytes(array<GoldilocksField, 4>(
state[0], state[1], state[2], state[3]
));
for (var i = 0u; i < 8u; i++) {
hash_le[i] = bswap32(second_output[7u - i]);
}
var below = cmp == 2u;
if (!below) {
for (var i = 0u; i < 8u; i++) {
let h = hash_le[7u - i];
let t = difficulty_target[7u - i];
if (h != t) {
below = h < t;
break;
}
}
}
if (below) {
// Try to claim the solution
if (atomicExchange(&results[0], 1u) == 0u) {
// We won! Write nonce and hash
// results layout: [0]=found, [1..16]=nonce, [17..32]=hash
for (var i = 0u; i < 16u; i++) {
atomicStore(&results[1u + i], current_nonce[i]);
crates/engine-gpu/src/kernels/mining_u64.wgsl
+573
View File
@@ -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
+650
View File
@@ -0,0 +1,650 @@
// 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),
// 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;
// Padded with a trailing zero so the internal layer can add the next round's
// constant unconditionally.
const RC_INTERNAL: array<u64, 23> = array<u64, 23>(
0x97f7798a784ad863lu, 0xd1d2bf082f60d4f0lu, 0x69a377a79f9ad206lu, 0xa9d06906a3858e24lu, 0x295275001eede5b5lu, 0x5874e441117bd746lu, 0x8a084bbba8ed86cclu, 0x3defd7645cde6425lu, 0x3998cfe6871cc137lu, 0x3e52ef8bca48314alu, 0x964a209f85dc9ecclu, 0x3fcc9ee82cc4577elu, 0x8e79b4a5d0096d6dlu, 0x8492362ad2392556lu, 0xee72f470262574d6lu, 0x1e0e18496da2444alu, 0x0f3a74bf215eaac6lu, 0x1b061b76a1c0ded3lu, 0x192c42d86803d7a6lu, 0xf6d49ff997ae0260lu, 0x3ec372e7a0fa3786lu, 0x5538cdf4f23445d3lu, 0lu
);
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);
}
// Sum accumulator: value = lo + carries * 2^64, folded once via 2^64 ≡ EPS64 (mod P).
struct Acc {
lo: u64,
carries: u32,
}
fn acc_add(a: Acc, b: u64) -> Acc {
let s = a.lo + b;
return Acc(s, a.carries + select(0u, 1u, s < a.lo));
}
fn acc_add2(a: Acc, b: Acc) -> Acc {
let s = a.lo + b.lo;
return Acc(s, a.carries + b.carries + select(0u, 1u, s < a.lo));
}
fn acc_fold(a: Acc) -> u64 {
let c = u64(a.carries);
let t = a.lo + ((c << 32u) - c);
return t + select(0lu, EPS64, t < a.lo);
}
struct U128 {
lo: u64,
hi: u64,
}
// Reduce a 128-bit value (lo + hi*2^64) mod P using
// 2^64 ≡ EPS64 and 2^96 ≡ -1 (mod P).
fn gf64_reduce(v: U128) -> u64 {
let hi_hi = v.hi >> 32u;
let hi_lo = v.hi & EPS64;
var t0 = v.lo - hi_hi;
t0 = t0 - select(0lu, EPS64, v.lo < hi_hi);
let t1 = hi_lo * EPS64;
let t2 = t0 + t1;
return t2 + select(0lu, EPS64, t2 < t0);
}
fn mul_wide(a: u64, b: u64) -> U128 {
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 = (ll >> 32u) + (lh & EPS64) + (hl & EPS64);
return U128((mid << 32u) | (ll & EPS64), hh + (lh >> 32u) + (hl >> 32u) + (mid >> 32u));
}
// (a*b + addend) mod P for b <= 2^64 - 2^32 (all MDS_DIAG entries): the addend's
// value and carries are folded into the 128-bit product before reduction.
fn gf64_mul_add(a: u64, b: u64, addend: Acc) -> u64 {
let v = mul_wide(a, b);
let lo = v.lo + addend.lo;
return gf64_reduce(U128(lo, v.hi + u64(addend.carries) + select(0lu, 1lu, lo < v.lo)));
}
fn gf64_mul(a: u64, b: u64) -> u64 {
return gf64_reduce(mul_wide(a, b));
}
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 = (ll >> 32u) + ((lh & EPS64) << 1u);
return gf64_reduce(U128((mid << 32u) | (ll & EPS64), hh + ((lh >> 32u) << 1u) + (mid >> 32u)));
}
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);
}
// 4x4 MDS circ(2, 3, 1, 1) on one chunk, results left unreduced.
fn mds4(x0: u64, x1: u64, x2: u64, x3: u64) -> array<Acc, 4> {
let t01 = acc_add(Acc(x0, 0u), x1);
let t23 = acc_add(Acc(x2, 0u), x3);
let t0123 = acc_add2(t01, t23);
let t01123 = acc_add(t0123, x1);
let t01233 = acc_add(t0123, x3);
return array<Acc, 4>(
acc_add2(t01123, t01),
acc_add2(t01123, acc_add(Acc(x2, 0u), x2)),
acc_add2(t01233, t23),
acc_add2(t01233, acc_add(Acc(x0, 0u), x0))
);
}
// External linear layer: 4x4 MDS on each chunk, then circulant sums, plus the
// next round's constants. Additions are accumulated unreduced (at most 27
// carries) and folded once per output.
fn ext_layer64(state: ptr<function, array<u64, 12>>, rc: array<u64, 12>) {
var y: array<Acc, 12>;
for (var chunk = 0u; chunk < 3u; chunk++) {
let o = chunk * 4u;
let m = mds4((*state)[o], (*state)[o + 1u], (*state)[o + 2u], (*state)[o + 3u]);
y[o] = m[0];
y[o + 1u] = m[1];
y[o + 2u] = m[2];
y[o + 3u] = m[3];
}
for (var k = 0u; k < 4u; k++) {
let s = acc_add2(acc_add2(y[k], y[k + 4u]), y[k + 8u]);
(*state)[k] = acc_fold(acc_add(acc_add2(y[k], s), rc[k]));
(*state)[k + 4u] = acc_fold(acc_add(acc_add2(y[k + 4u], s), rc[k + 4u]));
(*state)[k + 8u] = acc_fold(acc_add(acc_add2(y[k + 8u], s), rc[k + 8u]));
}
}
// Internal linear layer: diagonal matrix plus full sum, plus the next round's
// constant on element 0 (the S-box output, which is summed last).
fn int_layer64(state: ptr<function, array<u64, 12>>, rc0: u64) {
let s12 = acc_add(Acc((*state)[1], 0u), (*state)[2]);
let s34 = acc_add(Acc((*state)[3], 0u), (*state)[4]);
let s56 = acc_add(Acc((*state)[5], 0u), (*state)[6]);
let s78 = acc_add(Acc((*state)[7], 0u), (*state)[8]);
let s910 = acc_add(Acc((*state)[9], 0u), (*state)[10]);
let s1234 = acc_add2(s12, s34);
let s5678 = acc_add2(s56, s78);
let s91011 = acc_add(s910, (*state)[11]);
let sum = acc_add(acc_add2(acc_add2(s1234, s5678), s91011), (*state)[0]);
(*state)[0] = gf64_mul_add((*state)[0], MDS_DIAG[0], acc_add(sum, rc0));
(*state)[1] = gf64_mul_add((*state)[1], MDS_DIAG[1], sum);
(*state)[2] = gf64_mul_add((*state)[2], MDS_DIAG[2], sum);
(*state)[3] = gf64_mul_add((*state)[3], MDS_DIAG[3], sum);
(*state)[4] = gf64_mul_add((*state)[4], MDS_DIAG[4], sum);
(*state)[5] = gf64_mul_add((*state)[5], MDS_DIAG[5], sum);
(*state)[6] = gf64_mul_add((*state)[6], MDS_DIAG[6], sum);
(*state)[7] = gf64_mul_add((*state)[7], MDS_DIAG[7], sum);
(*state)[8] = gf64_mul_add((*state)[8], MDS_DIAG[8], sum);
(*state)[9] = gf64_mul_add((*state)[9], MDS_DIAG[9], sum);
(*state)[10] = gf64_mul_add((*state)[10], MDS_DIAG[10], sum);
(*state)[11] = gf64_mul_add((*state)[11], MDS_DIAG[11], sum);
}
fn sbox_lanes(state: ptr<function, array<u64, 12>>, lanes: u32) {
for (var i = 0u; i < lanes; i++) {
(*state)[i] = gf64_sbox((*state)[i]);
}
}
fn add_rc(state: ptr<function, array<u64, 12>>, rc: array<u64, 12>) {
for (var i = 0u; i < 12u; i++) {
(*state)[i] = gf64_add((*state)[i], rc[i]);
}
}
const RC_ZERO: array<u64, 12> = array<u64, 12>(0lu, 0lu, 0lu, 0lu, 0lu, 0lu, 0lu, 0lu, 0lu, 0lu, 0lu, 0lu);
// Constants added by each external layer, which feeds the S-box that follows it:
// entries 0..3 precede the initial rounds, entry 4 carries the first internal
// round's constant on element 0, entries 5..7 precede terminal rounds 1..3.
const RC_EXT: array<array<u64, 12>, 9> = array<array<u64, 12>, 9>(
RC_INITIAL[0], RC_INITIAL[1], RC_INITIAL[2], RC_INITIAL[3],
array<u64, 12>(RC_INTERNAL[0], 0lu, 0lu, 0lu, 0lu, 0lu, 0lu, 0lu, 0lu, 0lu, 0lu, 0lu),
RC_TERMINAL[1], RC_TERMINAL[2], RC_TERMINAL[3], RC_ZERO
);
// One loop drives all 30 rounds so each layer is emitted once: step 0 is the
// initial external layer, steps 1..4 and 27..30 are external rounds, steps
// 5..26 are internal rounds. Round constants are added by the linear layer
// preceding each S-box (the last internal layer is followed by the first
// terminal round's constants).
fn permute64(state: ptr<function, array<u64, 12>>) {
for (var k = 0u; k < 31u; k++) {
let is_ext = k < 5u || k > 26u;
if (k > 0u) {
sbox_lanes(state, select(1u, 12u, is_ext));
}
if (is_ext) {
ext_layer64(state, RC_EXT[select(k, k - 22u, k > 26u)]);
} else {
int_layer64(state, RC_INTERNAL[k - 4u]);
if (k == 26u) {
add_rc(state, RC_TERMINAL[0]);
}
}
}
}
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 = mid;
for (var i = 0u; i < 8u; i++) {
st[i] = gf64_add(st[i], u64(bswap32(current_nonce[7u - i])));
}
// Squeeze-and-compare phases share one inlined permutation: phase 0
// pads after absorbing, phase 1 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, and only candidates run phase 2 for the
// low half. Byte-swapped hash words are produced on demand.
var hash_le: array<u32, 16>;
var cmp = 0u;
var below = false;
for (var phase = 0u; phase < 3u; phase++) {
permute64(&st);
if (phase == 0u) {
st[0] = gf64_add(st[0], 1lu);
st[1] = gf64_add(st[1], 1lu);
continue;
}
var words: array<u32, 8>;
for (var i = 0u; i < 4u; i++) {
let c = gf64_canon(st[i]);
words[2u * i] = bswap32(u32(c & EPS64));
words[2u * i + 1u] = bswap32(u32(c >> 32u));
}
let base = select(15u, 7u, phase == 2u);
for (var i = 0u; i < 8u; i++) {
hash_le[base - i] = words[i];
}
if (phase == 1u) {
for (var i = 0u; i < 8u; i++) {
let h = words[i];
let t = tgt[15u - i];
if (h != t) {
cmp = select(2u, 1u, h > t);
break;
}
}
if (cmp == 1u) {
break;
}
} else {
below = cmp == 2u;
if (!below) {
for (var i = 0u; i < 8u; i++) {
let h = words[i];
let t = tgt[7u - i];
if (h != t) {
below = h < t;
break;
}
}
}
}
}
if (cmp == 1u) {
continue;
}
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, RC_ZERO);
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, 0lu);
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/mod.rs
+116
View File
@@ -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
+342 -74
View File
@@ -1,11 +1,14 @@
#![deny(rust_2018_idioms)]
#![forbid(unsafe_code)]
#![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;
@@ -27,7 +30,7 @@ struct GpuContext {
#[derive(Clone)]
struct GpuResources {
header_buffer: wgpu::Buffer,
midstate_buffer: wgpu::Buffer,
target_buffer: wgpu::Buffer,
start_nonce_buffer: wgpu::Buffer,
results_buffer: wgpu::Buffer,
@@ -37,26 +40,34 @@ struct GpuResources {
}
pub struct GpuEngine {
engine_id: usize,
contexts: Vec<Arc<GpuContext>>,
device_counter: AtomicUsize,
batch_size: u32,
throttle_ms: u64,
}
// Thread-local storage for consistent GPU device assignment per worker thread
static ENGINE_ID_COUNTER: AtomicUsize = AtomicUsize::new(0);
// Thread-local storage for consistent GPU device assignment per worker thread.
// Entries are tagged with the owning engine's id so resources created for one
// GpuEngine (and its devices) are never used with another (issue seen in benches
// where multiple engines exist in one process).
thread_local! {
static ASSIGNED_GPU_DEVICE: RefCell<Option<usize>> = const { RefCell::new(None) };
static WORKER_RESOURCES: RefCell<Option<GpuResources>> = const { RefCell::new(None) };
static ASSIGNED_GPU_DEVICE: RefCell<Option<(usize, usize)>> = const { RefCell::new(None) };
static WORKER_RESOURCES: RefCell<Option<(usize, GpuResources)>> = const { RefCell::new(None) };
/// Engine id whose GPU device was lost/unresponsive for this worker thread.
static DEVICE_LOST: RefCell<Option<usize>> = const { RefCell::new(None) };
}
impl GpuContext {
fn create_resources(&self) -> GpuResources {
let bind_group_layout = self.pipeline.get_bind_group_layout(0);
// Header: 8 u32s
let header_buffer = self.device.create_buffer(&wgpu::BufferDescriptor {
label: Some("Header Buffer"),
size: 32,
// Midstate: 12 felts as LE u32 pairs
let midstate_buffer = self.device.create_buffer(&wgpu::BufferDescriptor {
label: Some("Midstate Buffer"),
size: 96,
usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
mapped_at_creation: false,
});
@@ -113,7 +124,7 @@ impl GpuContext {
},
wgpu::BindGroupEntry {
binding: 1,
resource: header_buffer.as_entire_binding(),
resource: midstate_buffer.as_entire_binding(),
},
wgpu::BindGroupEntry {
binding: 2,
@@ -131,7 +142,7 @@ impl GpuContext {
});
GpuResources {
header_buffer,
midstate_buffer,
target_buffer,
start_nonce_buffer,
results_buffer,
@@ -142,6 +153,26 @@ 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 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 every kernel path).
#[allow(unsafe_code)]
fn create_trusted_shader(device: &wgpu::Device, shader_source: &str) -> wgpu::ShaderModule {
unsafe {
device.create_shader_module_trusted(
wgpu::ShaderModuleDescriptor {
label: Some("Mining Shader"),
source: wgpu::ShaderSource::Wgsl(shader_source.into()),
},
wgpu::ShaderRuntimeChecks::unchecked(),
)
}
}
/// Rank backends for mining: native compute APIs first.
fn backend_rank(backend: wgpu::Backend) -> u8 {
match backend {
@@ -160,6 +191,13 @@ fn backend_rank(backend: wgpu::Backend) -> u8 {
/// adapters are dropped and only the best-ranked backend present is kept.
/// Within a single backend each physical GPU appears exactly once, so rigs with
/// multiple identical cards keep every card.
///
/// When discrete GPUs are present, integrated GPUs (APUs) are skipped by default
/// to avoid resource contention and driver instability from mining on both
/// simultaneously. Set `allow_integrated` to true to override this behavior.
///
/// Note: This function no longer filters integrated GPUs - that decision is made
/// after initialization, so we can fall back to integrated if discrete fails.
fn select_adapters(infos: &[wgpu::AdapterInfo]) -> Vec<usize> {
let usable: Vec<usize> = (0..infos.len())
.filter(|&i| {
@@ -196,6 +234,8 @@ fn select_adapters(infos: &[wgpu::AdapterInfo]) -> Vec<usize> {
})
.collect();
// Sort discrete GPUs first, but keep all adapters for now
// (integrated filtering happens after init, based on what actually succeeded)
selected.sort_by_key(|&i| match infos[i].device_type {
wgpu::DeviceType::DiscreteGpu => 0,
wgpu::DeviceType::IntegratedGpu => 1,
@@ -204,15 +244,53 @@ fn select_adapters(infos: &[wgpu::AdapterInfo]) -> Vec<usize> {
selected
}
/// Filter initialized GPUs based on device types.
/// Returns indices of GPUs to keep.
///
/// Rules:
/// - If any discrete GPU initialized successfully and `allow_integrated` is false,
/// drop all integrated GPUs
/// - Otherwise keep all GPUs
///
/// This is extracted as a pure function for testability.
fn filter_initialized_gpus(
device_types: &[wgpu::DeviceType],
allow_integrated: bool,
) -> Vec<usize> {
let has_discrete = device_types.contains(&wgpu::DeviceType::DiscreteGpu);
if has_discrete && !allow_integrated {
// Keep only discrete GPUs
device_types
.iter()
.enumerate()
.filter(|(_, &dt)| dt != wgpu::DeviceType::IntegratedGpu)
.map(|(i, _)| i)
.collect()
} else {
// Keep all
(0..device_types.len()).collect()
}
}
impl GpuEngine {
/// Try to initialize the GPU engine with the given batch size and throttle (ms between batches).
///
/// # Arguments
/// * `batch_size` - Number of nonces per batch
/// * `throttle_ms` - Delay between batches in milliseconds (0 = no throttle)
/// * `allow_integrated` - If true, use integrated GPUs even when discrete GPUs are available
///
/// # Errors
///
/// Returns an error if:
/// - `batch_size` is zero (no work would be performed)
/// - No usable GPU adapters are found
pub fn try_new(batch_size: u32, throttle_ms: u64) -> Result<Self, Box<dyn std::error::Error>> {
pub fn try_new(
batch_size: u32,
throttle_ms: u64,
allow_integrated: bool,
) -> Result<Self, Box<dyn std::error::Error>> {
if batch_size == 0 {
return Err("batch_size must be non-zero".into());
}
@@ -221,17 +299,23 @@ impl GpuEngine {
match tokio::runtime::Handle::try_current() {
Ok(handle) => {
// We're inside a tokio runtime - use block_in_place to allow blocking
tokio::task::block_in_place(|| handle.block_on(Self::init(batch_size, throttle_ms)))
tokio::task::block_in_place(|| {
handle.block_on(Self::init(batch_size, throttle_ms, allow_integrated))
})
}
Err(_) => {
// No runtime exists - create a temporary one
let rt = tokio::runtime::Runtime::new()?;
rt.block_on(Self::init(batch_size, throttle_ms))
rt.block_on(Self::init(batch_size, throttle_ms, allow_integrated))
}
}
}
async fn init(batch_size: u32, throttle_ms: u64) -> Result<Self, Box<dyn std::error::Error>> {
async fn init(
batch_size: u32,
throttle_ms: u64,
allow_integrated: bool,
) -> Result<Self, Box<dyn std::error::Error>> {
log::info!(target: "gpu_engine", "Initializing WGPU...");
let instance = wgpu::Instance::new(&wgpu::InstanceDescriptor {
backends: wgpu::Backends::PRIMARY,
@@ -252,7 +336,14 @@ impl GpuEngine {
}
let mut adapters: Vec<Option<wgpu::Adapter>> = adapters.into_iter().map(Some).collect();
let mut contexts = Vec::new();
// Track successfully initialized contexts with their device type
struct InitializedGpu {
context: Arc<GpuContext>,
device_type: wgpu::DeviceType,
name: String,
}
let mut initialized: Vec<InitializedGpu> = Vec::new();
// Timeout for initializing each adapter (30 seconds should be plenty)
let init_timeout = std::time::Duration::from_secs(30);
@@ -269,11 +360,17 @@ impl GpuEngine {
);
log::debug!(target: "gpu_engine", "Adapter {i} raw info: {info:?}");
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: wgpu::Features::empty(),
required_features: if kernel.needs_int64() {
wgpu::Features::SHADER_INT64
} else {
wgpu::Features::empty()
},
required_limits: wgpu::Limits::default(),
memory_hints: Default::default(),
..Default::default()
@@ -316,11 +413,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 = include_str!("mining.wgsl");
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: Some("Mining Shader"),
source: wgpu::ShaderSource::Wgsl(shader_source.into()),
});
log::info!(
target: "gpu_engine",
"GPU device {i} ({}) using {} [{}]",
info.name,
kernel.label(),
kernel.id()
);
let shader = create_trusted_shader(&device, kernel.source());
let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
label: Some("Mining Pipeline"),
@@ -341,19 +441,44 @@ impl GpuEngine {
// Calculate vendor-specific configuration once during initialization
let optimal_workgroups = get_vendor_specific_dispatch(info, &device);
contexts.push(Arc::new(GpuContext {
device,
queue,
pipeline,
optimal_workgroups,
}));
initialized.push(InitializedGpu {
context: Arc::new(GpuContext {
device,
queue,
pipeline,
optimal_workgroups,
}),
device_type: info.device_type,
name: info.name.clone(),
});
}
if contexts.is_empty() {
if initialized.is_empty() {
log::error!(target: "gpu_engine", "No GPU adapters could be initialized successfully.");
return Err("No GPU adapters could be initialized".into());
}
// Filter integrated GPUs if discrete GPUs successfully initialized
let device_types: Vec<_> = initialized.iter().map(|g| g.device_type).collect();
let keep_indices = filter_initialized_gpus(&device_types, allow_integrated);
let contexts: Vec<Arc<GpuContext>> = initialized
.into_iter()
.enumerate()
.filter_map(|(i, g)| {
if keep_indices.contains(&i) {
Some(g.context)
} else {
log::info!(
target: "gpu_engine",
"Dropping integrated GPU (discrete GPU initialized successfully, use --allow-integrated to override): {}",
g.name
);
None
}
})
.collect();
log::info!(
target: "gpu_engine",
"GPU engine initialized with {} devices (batch size: {} nonces, throttle: {}ms)",
@@ -363,6 +488,7 @@ impl GpuEngine {
);
Ok(Self {
engine_id: ENGINE_ID_COUNTER.fetch_add(1, Ordering::SeqCst),
contexts,
device_counter: AtomicUsize::new(0),
batch_size,
@@ -376,11 +502,20 @@ impl GpuEngine {
}
/// Explicitly clear thread-local GPU resources.
/// Call this before thread exit to avoid TLS destruction order issues with wgpu.
/// Call this before dropping a `GpuEngine` (or before thread exit) so
/// buffers are destroyed while the device is still alive — otherwise the
/// next engine on this thread can reuse dead buffer IDs and panic with
/// `Buffer[…] does not exist`.
pub fn clear_worker_resources() {
WORKER_RESOURCES.with(|resources| {
*resources.borrow_mut() = None;
});
ASSIGNED_GPU_DEVICE.with(|assigned| {
*assigned.borrow_mut() = None;
});
DEVICE_LOST.with(|lost| {
*lost.borrow_mut() = None;
});
}
}
@@ -408,6 +543,13 @@ impl MinerEngine for GpuEngine {
return EngineStatus::Exhausted { hash_count: 0 };
}
// Check if this worker's GPU device was previously lost (for this engine)
let device_is_lost = DEVICE_LOST.with(|lost| *lost.borrow() == Some(self.engine_id));
if device_is_lost {
// Device was lost in a previous call - signal worker should exit
return EngineStatus::DeviceLost { hash_count: 0 };
}
// Empty or inverted range: nothing to do.
if range.start > range.end {
return EngineStatus::Exhausted { hash_count: 0 };
@@ -421,50 +563,42 @@ impl MinerEngine for GpuEngine {
// Use thread-local assignment for consistent worker-to-GPU mapping
let device_index = ASSIGNED_GPU_DEVICE.with(|assigned| {
let mut assigned_ref = assigned.borrow_mut();
if let Some(index) = *assigned_ref {
index
} else {
let index = if self.contexts.len() == 1 {
0
} else {
self.device_counter.fetch_add(1, Ordering::SeqCst) % self.contexts.len()
};
*assigned_ref = Some(index);
log::info!(
target: "gpu_engine",
"Worker thread assigned to GPU device {} (of {} total devices)",
index,
self.contexts.len()
);
index
match *assigned_ref {
Some((engine_id, index)) if engine_id == self.engine_id => index,
_ => {
let index = if self.contexts.len() == 1 {
0
} else {
self.device_counter.fetch_add(1, Ordering::SeqCst) % self.contexts.len()
};
*assigned_ref = Some((self.engine_id, index));
log::info!(
target: "gpu_engine",
"Worker thread assigned to GPU device {} (of {} total devices)",
index,
self.contexts.len()
);
index
}
}
});
let gpu_ctx = &self.contexts[device_index];
// Ensure resources are initialized for this thread
WORKER_RESOURCES.with(|resources_cell| {
// Ensure resources are initialized for this thread and belong to this engine
let resources = WORKER_RESOURCES.with(|resources_cell| {
let mut resources = resources_cell.borrow_mut();
if resources.is_none() {
*resources = Some(gpu_ctx.create_resources());
match &*resources {
Some((engine_id, res)) if *engine_id == self.engine_id => res.clone(),
_ => {
let res = gpu_ctx.create_resources();
*resources = Some((self.engine_id, res.clone()));
res
}
}
});
let resources = WORKER_RESOURCES
.with(|resources_cell| resources_cell.borrow().as_ref().unwrap().clone());
// Pre-convert header and target (only needs to be done once per job)
let mut header_u32s = [0u32; 8];
for (i, item) in header_u32s.iter_mut().enumerate() {
let chunk = &ctx.header[i * 4..(i + 1) * 4];
*item = u32::from_le_bytes([chunk[0], chunk[1], chunk[2], chunk[3]]);
}
gpu_ctx.queue.write_buffer(
&resources.header_buffer,
0,
bytemuck::cast_slice(&header_u32s),
);
// Pre-convert target (only needs to be done once per job)
let target_bytes = ctx.target.to_little_endian();
let mut target_u32s = [0u32; 16];
for i in 0..16 {
@@ -511,22 +645,27 @@ impl MinerEngine for GpuEngine {
};
}
// Calculate batch range
// Calculate batch range. The batch is additionally clamped so nonce
// increments never carry into the high 256 bits, which the midstate
// precompute (and the kernels) rely 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 batch_size_u512 = U512::from(self.batch_size);
let this_batch_size: u32 = if remaining > batch_size_u512 {
let this_batch_size: u32 = if cap > batch_size_u512 {
self.batch_size
} else {
// remaining fits in u32 since it's <= batch_size which is u32
remaining.low_u32()
// cap fits in u32 since it's <= batch_size which is u32
cap.low_u32()
};
// Run single batch
let batch_result =
run_single_batch(gpu_ctx, &resources, current_start, this_batch_size);
run_single_batch(gpu_ctx, &resources, ctx, current_start, this_batch_size);
match batch_result {
BatchResult::Found {
@@ -559,15 +698,18 @@ impl MinerEngine for GpuEngine {
total_hashes += hash_count;
}
BatchResult::DeviceLost => {
// GPU device is lost/unresponsive - log loudly and return cancelled
// This prevents spinning at 0 H/s indefinitely on a dead device
// 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);
log::error!(
target: "gpu_engine",
"GPU {} device lost or unresponsive - stopping worker. \
This GPU will not process further batches.",
device_index
);
return EngineStatus::Cancelled {
return EngineStatus::DeviceLost {
hash_count: total_hashes,
};
}
@@ -647,6 +789,7 @@ enum BatchResult {
fn run_single_batch(
gpu_ctx: &GpuContext,
resources: &GpuResources,
ctx: &JobContext,
batch_start: U512,
batch_size: u32,
) -> BatchResult {
@@ -679,6 +822,20 @@ fn run_single_batch(
.queue
.write_buffer(&resources.start_nonce_buffer, 0, &start_nonce_bytes);
// Precompute the sponge midstate for this batch (header + high nonce half)
let nonce_be = batch_start.to_big_endian();
let midstate = pow_core::mining_midstate(ctx.header, nonce_be[..32].try_into().unwrap());
let mut midstate_u32s = [0u32; 24];
for (i, felt) in midstate.iter().enumerate() {
midstate_u32s[2 * i] = *felt as u32;
midstate_u32s[2 * i + 1] = (*felt >> 32) as u32;
}
gpu_ctx.queue.write_buffer(
&resources.midstate_buffer,
0,
bytemuck::cast_slice(&midstate_u32s),
);
// Reset results buffer
const RESULTS_SIZE: usize = (1 + 16 + 16) * 4;
const ZEROS: [u8; RESULTS_SIZE] = [0; RESULTS_SIZE];
@@ -880,6 +1037,8 @@ mod adapter_selection_tests {
wgpu::Backend::Dx12,
),
];
// select_adapters returns all non-CPU adapters on best backend, discrete first
// Integrated filtering now happens after init in the init() function
assert_eq!(select_adapters(&infos), vec![1, 0]);
}
@@ -906,14 +1065,33 @@ mod adapter_selection_tests {
}
#[test]
fn dx12_only_machine_keeps_all_dx12_adapters() {
fn dx12_only_machine_returns_all_adapters_sorted() {
let infos = [
info("iGPU", wgpu::DeviceType::IntegratedGpu, wgpu::Backend::Dx12),
info("dGPU", wgpu::DeviceType::DiscreteGpu, wgpu::Backend::Dx12),
];
// select_adapters returns both, discrete first (integrated filtering is post-init)
assert_eq!(select_adapters(&infos), vec![1, 0]);
}
#[test]
fn integrated_only_machine_keeps_integrated() {
// When no discrete GPU exists, integrated GPUs should be used
let infos = [
info(
"Intel UHD",
wgpu::DeviceType::IntegratedGpu,
wgpu::Backend::Vulkan,
),
info(
"AMD Vega 8",
wgpu::DeviceType::IntegratedGpu,
wgpu::Backend::Vulkan,
),
];
assert_eq!(select_adapters(&infos), vec![0, 1]);
}
#[test]
fn software_only_environment_selects_nothing() {
let infos = [info(
@@ -928,4 +1106,94 @@ mod adapter_selection_tests {
fn empty_enumeration_selects_nothing() {
assert!(select_adapters(&[]).is_empty());
}
/// Exact scenario from Windows ASUS laptop with RX 560X + Vega 8 APU.
/// Both GPUs appear on both Vulkan and Dx12 backends.
/// select_adapters returns both Vulkan adapters (discrete first).
/// The init() function will later drop the integrated one if discrete succeeds.
#[test]
fn windows_amd_discrete_plus_apu_selects_both_on_best_backend() {
let infos = [
info(
"Microsoft Basic Render Driver",
wgpu::DeviceType::Cpu,
wgpu::Backend::Dx12,
),
info(
"AMD Radeon(TM) Vega 8 Graphics",
wgpu::DeviceType::IntegratedGpu,
wgpu::Backend::Dx12,
),
info(
"Radeon RX 560X",
wgpu::DeviceType::DiscreteGpu,
wgpu::Backend::Dx12,
),
info(
"AMD Radeon(TM) Vega 8 Graphics",
wgpu::DeviceType::IntegratedGpu,
wgpu::Backend::Vulkan,
),
info(
"Radeon RX 560X",
wgpu::DeviceType::DiscreteGpu,
wgpu::Backend::Vulkan,
),
];
// select_adapters returns both Vulkan adapters, discrete first
// - Index 0: Skipped (CPU emulated)
// - Index 1, 2: Skipped (Dx12 lower priority than Vulkan)
// - Index 4: Selected first (discrete, Vulkan)
// - Index 3: Selected second (integrated, Vulkan)
assert_eq!(select_adapters(&infos), vec![4, 3]);
}
// Tests for filter_initialized_gpus (post-init filtering)
#[test]
fn filter_discrete_present_drops_integrated() {
use wgpu::DeviceType::*;
// Discrete at index 0, integrated at index 1
let types = vec![DiscreteGpu, IntegratedGpu];
assert_eq!(filter_initialized_gpus(&types, false), vec![0]);
}
#[test]
fn filter_discrete_failed_keeps_integrated() {
use wgpu::DeviceType::*;
// Only integrated initialized (discrete failed/timed out)
let types = vec![IntegratedGpu];
assert_eq!(filter_initialized_gpus(&types, false), vec![0]);
}
#[test]
fn filter_allow_integrated_keeps_both() {
use wgpu::DeviceType::*;
let types = vec![DiscreteGpu, IntegratedGpu];
// With allow_integrated=true, keep both
assert_eq!(filter_initialized_gpus(&types, true), vec![0, 1]);
}
#[test]
fn filter_multiple_discrete_keeps_all_discrete() {
use wgpu::DeviceType::*;
let types = vec![DiscreteGpu, DiscreteGpu, IntegratedGpu];
// Drops integrated, keeps both discrete
assert_eq!(filter_initialized_gpus(&types, false), vec![0, 1]);
}
#[test]
fn filter_only_discrete_keeps_all() {
use wgpu::DeviceType::*;
let types = vec![DiscreteGpu, DiscreteGpu];
assert_eq!(filter_initialized_gpus(&types, false), vec![0, 1]);
}
#[test]
fn filter_multiple_integrated_no_discrete_keeps_all() {
use wgpu::DeviceType::*;
let types = vec![IntegratedGpu, IntegratedGpu];
// No discrete, so keep all integrated
assert_eq!(filter_initialized_gpus(&types, false), vec![0, 1]);
}
}
crates/engine-gpu/src/main.rs
+107 -151
View File
@@ -3,174 +3,130 @@ mod tests;
#[tokio::main]
async fn main() -> Result<(), Box<dyn std::error::Error>> {
run().await
}
async fn run() -> Result<(), Box<dyn std::error::Error>> {
// Setup GPU device and queue
let instance = wgpu::Instance::new(&wgpu::InstanceDescriptor {
backends: wgpu::Backends::METAL, // Force Metal on Apple
backends: wgpu::Backends::PRIMARY,
..Default::default()
});
let adapter = instance
.request_adapter(&wgpu::RequestAdapterOptions::default())
.await
.unwrap();
let (device, queue) = adapter
.request_device(&wgpu::DeviceDescriptor::default())
.await?;
.expect("no GPU adapter");
println!("Running Poseidon2 GPU Component Tests...\n");
let mut failures = 0usize;
let has_int64 = adapter.features().contains(wgpu::Features::SHADER_INT64);
// Run all component tests
if let Err(e) = tests::test_gf_from_const(&device, &queue).await {
eprintln!("❌ gf_from_const tests failed: {}", e);
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(kernel.label()),
required_features: if kernel.needs_int64() {
wgpu::Features::SHADER_INT64
} else {
wgpu::Features::empty()
},
..Default::default()
})
.await?;
failures += run_suite(&device, &queue, kernel.source(), kernel.label()).await;
}
if let Err(e) = tests::test_gf_mul(&device, &queue).await {
eprintln!("❌ gf_mul tests failed: {}", e);
if failures > 0 {
return Err(format!("{failures} test group(s) failed").into());
}
if let Err(e) = tests::test_sbox(&device, &queue).await {
eprintln!("❌ S-box tests failed: {}", e);
}
if let Err(e) = tests::test_mds_matrix(&device, &queue).await {
eprintln!("❌ MDS matrix tests failed: {}", e);
}
if let Err(e) = tests::test_internal_linear_layer(&device, &queue).await {
eprintln!("❌ Internal linear layer tests failed: {}", e);
}
if let Err(e) = tests::test_external_linear_layer(&device, &queue).await {
eprintln!("❌ External linear layer tests failed: {}", e);
}
if let Err(e) = tests::test_poseidon2_initial_external_rounds(&device, &queue).await {
eprintln!("❌ Initial external rounds tests failed: {}", e);
}
if let Err(e) = tests::test_poseidon2_terminal_external_rounds(&device, &queue).await {
eprintln!("❌ Terminal external rounds tests failed: {}", e);
}
if let Err(e) = tests::test_poseidon2_constants_verification(&device, &queue).await {
eprintln!("❌ Constants verification test failed: {}", e);
}
if let Err(e) = tests::test_poseidon2_internal_constants_verification(&device, &queue).await {
eprintln!("❌ Internal constants verification test failed: {}", e);
}
if let Err(e) = tests::test_poseidon2_internal_rounds_only(&device, &queue).await {
eprintln!("❌ Internal rounds only test failed: {}", e);
}
if let Err(e) =
tests::test_poseidon2_terminal_external_constants_verification(&device, &queue).await
{
eprintln!(
"❌ Terminal external constants verification test failed: {}",
e
);
}
if let Err(e) = tests::test_poseidon2_permutation(&device, &queue).await {
eprintln!("❌ Poseidon2 permutation tests failed: {}", e);
}
if let Err(e) = tests::test_bytes_to_field_elements(&device, &queue).await {
eprintln!("❌ Bytes to field elements tests failed: {}", e);
}
if let Err(e) = tests::test_field_elements_to_bytes(&device, &queue).await {
eprintln!("❌ Field elements to bytes tests failed: {}", e);
}
if let Err(e) = tests::test_poseidon2_squeeze_twice(&device, &queue).await {
eprintln!("❌ Poseidon2 squeeze-twice tests failed: {}", e);
}
if let Err(e) = tests::test_hash_squeeze_twice(&device, &queue).await {
eprintln!("❌ Hash squeeze twice tests failed: {}", e);
}
if let Err(e) = end_to_end_tests::test_end_to_end_mining(&device, &queue).await {
eprintln!("❌ End-to-end mining test failed: {}", e);
}
println!("\nAll tests completed!");
if let Err(e) = end_to_end_tests::test_end_to_end_mining(&device, &queue).await {
eprintln!("❌ End-to-end mining test failed: {}", e);
}
println!("\nAll tests completed!");
// generate_correct_wgsl_constants();
Ok(())
}
#[allow(dead_code)]
fn generate_correct_wgsl_constants() {
use qp_poseidon_constants::*;
println!("🔧 Generating correct WGSL constants...");
println!("// Initial external round constants (4 rounds x 12 elements)");
println!("const INITIAL_EXTERNAL_CONSTANTS: array<array<array<u32, 2>, 12>, 4> = array<array<array<u32, 2>, 12>, 4>(");
for (round_idx, round) in POSEIDON2_INITIAL_EXTERNAL_CONSTANTS_RAW.iter().enumerate() {
println!(" array<array<u32, 2>, 12>(");
for (elem_idx, &value) in round.iter().enumerate() {
let low = value as u32;
let high = (value >> 32) as u32;
if elem_idx == 11 {
println!(" array<u32, 2>({}u, {}u)", low, high);
} else {
println!(" array<u32, 2>({}u, {}u),", low, high);
async fn run_suite(
device: &wgpu::Device,
queue: &wgpu::Queue,
shader_src: &str,
label: &str,
) -> usize {
println!("\n==== Running Poseidon2 GPU tests against {label} ====\n");
let mut failures = 0usize;
macro_rules! run {
($name:literal, $fut:expr) => {
if let Err(e) = $fut.await {
eprintln!("❌ {} failed: {}", $name, e);
failures += 1;
}
}
if round_idx == 3 {
println!(" )");
} else {
println!(" ),");
}
};
}
println!(");");
println!("\n// Terminal external round constants (4 rounds x 12 elements)");
println!("const TERMINAL_EXTERNAL_CONSTANTS: array<array<array<u32, 2>, 12>, 4> = array<array<array<u32, 2>, 12>, 4>(");
run!(
"gf_from_const",
tests::test_gf_from_const(device, queue, shader_src)
);
run!("gf_mul", tests::test_gf_mul(device, queue, shader_src));
run!("sbox", tests::test_sbox(device, queue, shader_src));
run!(
"mds matrix",
tests::test_mds_matrix(device, queue, shader_src)
);
run!(
"internal linear layer",
tests::test_internal_linear_layer(device, queue, shader_src)
);
run!(
"external linear layer",
tests::test_external_linear_layer(device, queue, shader_src)
);
run!(
"initial external rounds",
tests::test_poseidon2_initial_external_rounds(device, queue, shader_src)
);
run!(
"terminal external rounds",
tests::test_poseidon2_terminal_external_rounds(device, queue, shader_src)
);
run!(
"constants verification",
tests::test_poseidon2_constants_verification(device, queue, shader_src)
);
run!(
"internal constants verification",
tests::test_poseidon2_internal_constants_verification(device, queue, shader_src)
);
run!(
"internal rounds only",
tests::test_poseidon2_internal_rounds_only(device, queue, shader_src)
);
run!(
"terminal external constants verification",
tests::test_poseidon2_terminal_external_constants_verification(device, queue, shader_src)
);
run!(
"poseidon2 permutation",
tests::test_poseidon2_permutation(device, queue, shader_src)
);
run!(
"bytes to field elements",
tests::test_bytes_to_field_elements(device, queue, shader_src)
);
run!(
"field elements to bytes",
tests::test_field_elements_to_bytes(device, queue, shader_src)
);
run!(
"poseidon2 squeeze-twice",
tests::test_poseidon2_squeeze_twice(device, queue, shader_src)
);
run!(
"hash squeeze twice",
tests::test_hash_squeeze_twice(device, queue, shader_src)
);
run!(
"end-to-end mining",
end_to_end_tests::test_end_to_end_mining(device, queue, shader_src)
);
for (round_idx, round) in POSEIDON2_TERMINAL_EXTERNAL_CONSTANTS_RAW.iter().enumerate() {
println!(" array<array<u32, 2>, 12>(");
for (elem_idx, &value) in round.iter().enumerate() {
let low = value as u32;
let high = (value >> 32) as u32;
if elem_idx == 11 {
println!(" array<u32, 2>({}u, {}u)", low, high);
} else {
println!(" array<u32, 2>({}u, {}u),", low, high);
}
}
if round_idx == 3 {
println!(" )");
} else {
println!(" ),");
}
}
println!(");");
println!("\n// Internal round constants (22 values)");
println!("const INTERNAL_CONSTANTS: array<array<u32, 2>, 22> = array<array<u32, 2>, 22>(");
for (idx, &value) in POSEIDON2_INTERNAL_CONSTANTS_RAW.iter().enumerate() {
let low = value as u32;
let high = (value >> 32) as u32;
if idx == 21 {
println!(" array<u32, 2>({}u, {}u)", low, high);
} else {
println!(" array<u32, 2>({}u, {}u),", low, high);
}
}
println!(");");
failures
}
crates/engine-gpu/src/tests.rs
+36 -18
View File
@@ -1166,6 +1166,7 @@ fn generate_true_internal_only_test_vectors() -> Vec<InternalRoundsTestCase> {
pub async fn test_poseidon2_permutation(
device: &wgpu::Device,
queue: &wgpu::Queue,
shader_src: &str,
) -> Result<(), Box<dyn std::error::Error>> {
let test_vectors = generate_poseidon2_test_vectors();
let total_tests = test_vectors.len();
@@ -1202,7 +1203,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
}}
}}
",
include_str!("mining.wgsl")
shader_src
);
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -1370,6 +1371,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
pub async fn test_poseidon2_initial_external_rounds(
device: &wgpu::Device,
queue: &wgpu::Queue,
shader_src: &str,
) -> Result<(), Box<dyn std::error::Error>> {
// Test just the initial 4 external rounds
let test_vectors = generate_gf_from_const_test_vectors();
@@ -1415,7 +1417,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
}}
}}
",
include_str!("mining.wgsl")
shader_src
);
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -1630,6 +1632,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
pub async fn test_poseidon2_terminal_external_rounds(
device: &wgpu::Device,
queue: &wgpu::Queue,
shader_src: &str,
) -> Result<(), Box<dyn std::error::Error>> {
// Test just the terminal 4 external rounds
let _total_tests = 2;
@@ -1674,7 +1677,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
}}
}}
",
include_str!("mining.wgsl")
shader_src
);
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -1888,6 +1891,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
pub async fn test_gf_from_const(
device: &wgpu::Device,
queue: &wgpu::Queue,
shader_src: &str,
) -> Result<(), Box<dyn std::error::Error>> {
// Test the gf_from_const function that converts constant arrays to GoldilocksField
let test_vectors = generate_gf_from_const_test_vectors();
@@ -1908,7 +1912,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
output_fields[index] = gf_from_const(input_constants[index]);
}}
",
include_str!("mining.wgsl")
shader_src
);
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -2023,6 +2027,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
pub async fn test_poseidon2_constants_verification(
device: &wgpu::Device,
queue: &wgpu::Queue,
shader_src: &str,
) -> Result<(), Box<dyn std::error::Error>> {
let mut passed_tests = 0;
let mut failed_tests = Vec::new();
@@ -2049,7 +2054,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
}}
}}
",
include_str!("mining.wgsl")
shader_src
);
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -2160,6 +2165,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
pub async fn test_poseidon2_internal_rounds_only(
device: &wgpu::Device,
queue: &wgpu::Queue,
shader_src: &str,
) -> Result<(), Box<dyn std::error::Error>> {
// Test ONLY the internal rounds (22 rounds) without initial/terminal external rounds
let test_vectors = generate_true_internal_only_test_vectors();
@@ -2209,7 +2215,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
}}
}}
",
include_str!("mining.wgsl")
shader_src
);
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -2343,6 +2349,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
pub async fn test_poseidon2_internal_constants_verification(
device: &wgpu::Device,
queue: &wgpu::Queue,
shader_src: &str,
) -> Result<(), Box<dyn std::error::Error>> {
let mut passed_tests = 0;
let mut failed_tests = Vec::new();
@@ -2364,7 +2371,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
}}
}}
",
include_str!("mining.wgsl")
shader_src
);
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -2462,6 +2469,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
pub async fn test_poseidon2_terminal_external_rounds_issue(
device: &wgpu::Device,
queue: &wgpu::Queue,
shader_src: &str,
) -> Result<(), Box<dyn std::error::Error>> {
println!("🔍 Testing terminal external rounds issue for sequential input...");
@@ -2502,7 +2510,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
}}
}}
",
include_str!("mining.wgsl")
shader_src
);
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -2630,6 +2638,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
pub async fn test_poseidon2_terminal_external_constants_verification(
device: &wgpu::Device,
queue: &wgpu::Queue,
shader_src: &str,
) -> Result<(), Box<dyn std::error::Error>> {
let mut passed_tests = 0;
let mut failed_tests = Vec::new();
@@ -2656,7 +2665,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
}}
}}
",
include_str!("mining.wgsl")
shader_src
);
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -2764,6 +2773,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
pub async fn test_gf_mul(
device: &wgpu::Device,
queue: &wgpu::Queue,
shader_src: &str,
) -> Result<(), Box<dyn std::error::Error>> {
let test_vectors = generate_gf_mul_test_vectors();
let total_tests = test_vectors.len();
@@ -2786,7 +2796,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
output[i] = gf_mul(input_a[i], input_b[i]);
}}
",
include_str!("mining.wgsl")
shader_src
);
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -2947,6 +2957,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
pub async fn test_mds_matrix(
device: &wgpu::Device,
queue: &wgpu::Queue,
shader_src: &str,
) -> Result<(), Box<dyn std::error::Error>> {
let test_vectors = generate_mds_test_vectors();
let total_tests = test_vectors.len();
@@ -2955,7 +2966,7 @@ pub async fn test_mds_matrix(
let mut failure_details = Vec::new();
// Read the mining shader source and create a test wrapper
let mining_shader_source = include_str!("mining.wgsl");
let mining_shader_source = shader_src;
let shader_source = format!(
"
{}
@@ -3215,6 +3226,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
pub async fn test_internal_linear_layer(
device: &wgpu::Device,
queue: &wgpu::Queue,
shader_src: &str,
) -> Result<(), Box<dyn std::error::Error>> {
let test_vectors = generate_internal_linear_layer_test_vectors();
let total_tests = test_vectors.len();
@@ -3223,7 +3235,7 @@ pub async fn test_internal_linear_layer(
let mut failure_details = Vec::new();
// Read the mining shader source and create a test wrapper
let mining_shader_source = include_str!("mining.wgsl");
let mining_shader_source = shader_src;
let shader_source = format!(
"
{}
@@ -3448,6 +3460,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
pub async fn test_external_linear_layer(
device: &wgpu::Device,
queue: &wgpu::Queue,
shader_src: &str,
) -> Result<(), Box<dyn std::error::Error>> {
let test_vectors = generate_external_linear_layer_test_vectors();
let total_tests = test_vectors.len();
@@ -3485,7 +3498,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
}}
}}
",
include_str!("mining.wgsl")
shader_src
);
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -3625,6 +3638,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
pub async fn test_sbox(
device: &wgpu::Device,
queue: &wgpu::Queue,
shader_src: &str,
) -> Result<(), Box<dyn std::error::Error>> {
let test_vectors = generate_sbox_test_vectors();
let total_tests = test_vectors.len();
@@ -3662,7 +3676,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
output_data[index] = result;
}}
",
include_str!("mining.wgsl"),
shader_src,
batch.len()
);
@@ -3807,6 +3821,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
pub async fn test_poseidon2_squeeze_twice(
device: &wgpu::Device,
queue: &wgpu::Queue,
shader_src: &str,
) -> Result<(), Box<dyn std::error::Error>> {
let test_vectors = generate_poseidon2_squeeze_twice_test_vectors();
let total_tests = test_vectors.len();
@@ -3842,7 +3857,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
}}
}}
",
include_str!("mining.wgsl")
shader_src
);
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -3985,6 +4000,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
pub async fn test_bytes_to_field_elements(
device: &wgpu::Device,
queue: &wgpu::Queue,
shader_src: &str,
) -> Result<(), Box<dyn std::error::Error>> {
let test_vectors = generate_bytes_to_field_test_vectors();
let total_tests = test_vectors.len();
@@ -4020,7 +4036,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
}}
}}
",
include_str!("mining.wgsl")
shader_src
);
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -4168,6 +4184,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
pub async fn test_field_elements_to_bytes(
device: &wgpu::Device,
queue: &wgpu::Queue,
shader_src: &str,
) -> Result<(), Box<dyn std::error::Error>> {
let test_vectors = generate_field_to_bytes_test_vectors();
let total_tests = test_vectors.len();
@@ -4203,7 +4220,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
}}
}}
",
include_str!("mining.wgsl")
shader_src
);
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -4336,6 +4353,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
pub async fn test_hash_squeeze_twice(
device: &wgpu::Device,
queue: &wgpu::Queue,
shader_src: &str,
) -> Result<(), Box<dyn std::error::Error>> {
let test_vectors = generate_hash_squeeze_twice_test_vectors();
let total_tests = test_vectors.len();
@@ -4371,7 +4389,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {{
}}
}}
",
include_str!("mining.wgsl")
shader_src
);
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
crates/miner-cli/Cargo.toml
+1
View File
@@ -8,6 +8,7 @@ description = "CLI binary to run the Quantus External Miner service"
[dependencies]
miner-service = { path = "../miner-service" }
metrics = { path = "../metrics", features = ["http-exporter"] }
quic-transport = { path = "../quic-transport" }
clap = { workspace = true, features = ["derive", "env"] }
tokio = { workspace = true, features = ["full"] }
env_logger = { workspace = true }
crates/miner-cli/src/main.rs
+173 -25
View File
@@ -3,6 +3,7 @@ use engine_cpu::{AtomicBoolCancelCheck, EngineRange, MinerEngine};
use miner_service::{run, ServiceConfig};
use primitive_types::U512;
use rand::RngCore;
use std::path::PathBuf;
use std::sync::atomic::AtomicBool;
use std::sync::Arc;
use std::thread;
@@ -20,6 +21,34 @@ enum Command {
#[arg(long, env = "MINER_NODE_ADDR", default_value = "127.0.0.1:9833")]
node_addr: std::net::SocketAddr,
/// Shared auth token from the node's `miner-auth-token` file
/// (`<base-path>/chains/<chain>/miner-auth-token`). Prefer `--auth-token-file`.
#[arg(long, env = "MINER_AUTH_TOKEN", conflicts_with = "auth_token_file")]
auth_token: Option<String>,
/// Path to the node's `miner-auth-token` file (trimmed). Preferred over
/// `--auth-token` so the secret is not placed on the command line.
#[arg(long, env = "MINER_AUTH_TOKEN_FILE", conflicts_with = "auth_token")]
auth_token_file: Option<PathBuf>,
/// SHA-256 fingerprint of the node's miner TLS certificate (64 hex chars).
/// Prefer `--tls-cert-sha256-file` pointing at `miner-tls-cert-sha256`
/// (the node also logs this fingerprint).
#[arg(
long,
env = "MINER_TLS_CERT_SHA256",
conflicts_with = "tls_cert_sha256_file"
)]
tls_cert_sha256: Option<String>,
/// Path to the node's `miner-tls-cert-sha256` file.
#[arg(
long,
env = "MINER_TLS_CERT_SHA256_FILE",
conflicts_with = "tls_cert_sha256"
)]
tls_cert_sha256_file: Option<PathBuf>,
/// Number of CPU worker threads to use for mining (default: auto-detect)
#[arg(long = "cpu-workers", env = "MINER_CPU_WORKERS")]
cpu_workers: Option<usize>,
@@ -29,11 +58,11 @@ enum Command {
gpu_devices: Option<usize>,
/// GPU batch size in nonces - controls how often GPU checks for cancellation
#[arg(long = "gpu-batch-size", env = "MINER_GPU_BATCH_SIZE", default_value_t = DEFAULT_GPU_BATCH_SIZE)]
#[arg(long = "gpu-batch-size", env = "MINER_GPU_BATCH_SIZE", default_value_t = DEFAULT_GPU_BATCH_SIZE, value_parser = clap::value_parser!(u32).range(1..))]
gpu_batch_size: u32,
/// CPU batch size in hashes - controls how often CPU checks for cancellation
#[arg(long = "cpu-batch-size", env = "MINER_CPU_BATCH_SIZE", default_value_t = DEFAULT_CPU_BATCH_SIZE)]
#[arg(long = "cpu-batch-size", env = "MINER_CPU_BATCH_SIZE", default_value_t = DEFAULT_CPU_BATCH_SIZE, value_parser = clap::value_parser!(u64).range(1..))]
cpu_batch_size: u64,
/// Port for Prometheus metrics HTTP endpoint (default: 9900)
@@ -52,6 +81,12 @@ enum Command {
)]
gpu_throttle_ms: u64,
/// Allow integrated GPUs (APUs) even when discrete GPUs are available.
/// By default, integrated GPUs are skipped when a discrete GPU is present
/// to avoid resource contention and driver instability.
#[arg(long = "allow-integrated", env = "MINER_ALLOW_INTEGRATED")]
allow_integrated: bool,
/// Enable verbose logging
#[arg(short, long, env = "MINER_VERBOSE")]
verbose: bool,
@@ -68,17 +103,21 @@ enum Command {
gpu_devices: Option<usize>,
/// GPU batch size in nonces - controls how often GPU checks for cancellation
#[arg(long = "gpu-batch-size", env = "MINER_GPU_BATCH_SIZE", default_value_t = DEFAULT_GPU_BATCH_SIZE)]
#[arg(long = "gpu-batch-size", env = "MINER_GPU_BATCH_SIZE", default_value_t = DEFAULT_GPU_BATCH_SIZE, value_parser = clap::value_parser!(u32).range(1..))]
gpu_batch_size: u32,
/// CPU batch size in hashes - controls how often CPU checks for cancellation
#[arg(long = "cpu-batch-size", env = "MINER_CPU_BATCH_SIZE", default_value_t = DEFAULT_CPU_BATCH_SIZE)]
#[arg(long = "cpu-batch-size", env = "MINER_CPU_BATCH_SIZE", default_value_t = DEFAULT_CPU_BATCH_SIZE, value_parser = clap::value_parser!(u64).range(1..))]
cpu_batch_size: u64,
/// Benchmark duration in seconds (default: 10)
#[arg(short, long, default_value_t = 10)]
duration: u64,
/// Allow integrated GPUs (APUs) even when discrete GPUs are available
#[arg(long = "allow-integrated", env = "MINER_ALLOW_INTEGRATED")]
allow_integrated: bool,
/// Enable verbose logging
#[arg(short, long, env = "MINER_VERBOSE")]
verbose: bool,
@@ -99,23 +138,53 @@ async fn main() {
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).");
eprintln!("Example: quantus-miner serve --node-addr 127.0.0.1:9833");
eprintln!(
"Example: quantus-miner serve --node-addr 127.0.0.1:9833 \
--auth-token-file /path/to/miner-auth-token \
--tls-cert-sha256-file /path/to/miner-tls-cert-sha256"
);
std::process::exit(1);
};
match command {
Command::Serve {
node_addr,
auth_token,
auth_token_file,
tls_cert_sha256,
tls_cert_sha256_file,
cpu_workers,
gpu_devices,
gpu_batch_size,
cpu_batch_size,
gpu_throttle_ms,
metrics_port,
allow_integrated,
verbose,
} => {
init_logger(verbose);
let auth_token = match resolve_auth_token(auth_token, auth_token_file) {
Ok(token) => token,
Err(e) => {
eprintln!("Error: {e}");
std::process::exit(1);
}
};
let tls_cert_sha256 =
match resolve_tls_cert_sha256(tls_cert_sha256, tls_cert_sha256_file) {
Ok(fp) => fp,
Err(e) => {
eprintln!("Error: {e}");
std::process::exit(1);
}
};
// Fail closed on permanent misconfig before metrics/workers start.
if let Err(e) = quic_transport::validate_auth_config(&auth_token, &tls_cert_sha256) {
eprintln!("Error: {e}");
std::process::exit(1);
}
log::info!("Starting external miner service...");
// Start metrics HTTP server
@@ -130,11 +199,14 @@ async fn main() {
let config = ServiceConfig {
node_addr,
auth_token,
tls_cert_sha256,
cpu_workers,
gpu_devices,
gpu_batch_size,
cpu_batch_size,
gpu_throttle_ms,
allow_integrated,
};
if let Err(e) = run(config).await {
@@ -149,6 +221,7 @@ async fn main() {
gpu_batch_size,
cpu_batch_size,
duration,
allow_integrated,
verbose,
} => {
init_logger(verbose);
@@ -158,12 +231,64 @@ async fn main() {
gpu_batch_size,
cpu_batch_size,
duration,
allow_integrated,
)
.await;
}
}
}
fn resolve_auth_token(
auth_token: Option<String>,
auth_token_file: Option<PathBuf>,
) -> Result<String, String> {
resolve_required_secret(
auth_token,
auth_token_file,
"--auth-token",
"--auth-token-file",
"miner auth token required: pass --auth-token-file <PATH> to the node's \
miner-auth-token file (or --auth-token <TOKEN>)",
)
}
fn resolve_tls_cert_sha256(value: Option<String>, file: Option<PathBuf>) -> Result<String, String> {
resolve_required_secret(
value,
file,
"--tls-cert-sha256",
"--tls-cert-sha256-file",
"TLS cert fingerprint required: pass --tls-cert-sha256-file <PATH> to the node's \
miner-tls-cert-sha256 file (or --tls-cert-sha256 <HEX>; also printed in node logs)",
)
}
fn resolve_required_secret(
value: Option<String>,
file: Option<PathBuf>,
value_flag: &str,
file_flag: &str,
missing_msg: &str,
) -> Result<String, String> {
if let Some(value) = value {
let value = value.trim().to_string();
if value.is_empty() {
return Err(format!("{value_flag} is empty"));
}
return Ok(value);
}
if let Some(path) = file {
let contents = std::fs::read_to_string(&path)
.map_err(|e| format!("failed to read {file_flag} {}: {}", path.display(), e))?;
let value = contents.trim().to_string();
if value.is_empty() {
return Err(format!("{file_flag} {} is empty", path.display()));
}
return Ok(value);
}
Err(missing_msg.into())
}
fn init_logger(verbose: bool) {
if std::env::var("RUST_LOG").is_err() {
// Filter out noisy wgpu/naga shader compilation logs
@@ -183,18 +308,23 @@ async fn run_benchmark(
gpu_batch_size: u32,
cpu_batch_size: u64,
duration: u64,
allow_integrated: bool,
) {
let effective_cpu_workers = cpu_workers.unwrap_or_else(num_cpus::get);
// Initialize GPU engine (no throttle for benchmark)
let (gpu_engine, effective_gpu_devices) =
match miner_service::resolve_gpu_configuration(gpu_devices, gpu_batch_size, 0) {
Ok((engine, count)) => (engine, count),
Err(e) => {
eprintln!("❌ ERROR: {}", e);
std::process::exit(1);
}
};
let (gpu_engine, effective_gpu_devices) = match miner_service::resolve_gpu_configuration(
gpu_devices,
gpu_batch_size,
0,
allow_integrated,
) {
Ok((engine, count)) => (engine, count),
Err(e) => {
eprintln!("❌ ERROR: {}", e);
std::process::exit(1);
}
};
let total_workers = effective_cpu_workers + effective_gpu_devices;
@@ -206,7 +336,13 @@ async fn run_benchmark(
num_cpus::get()
);
println!("GPU Devices: {}", effective_gpu_devices);
println!("Duration: {} seconds", duration);
if effective_cpu_workers > 0 {
println!("CPU batch size: {cpu_batch_size} hashes");
}
if effective_gpu_devices > 0 {
println!("GPU batch size: {gpu_batch_size} nonces");
}
println!("Duration: {duration} seconds");
println!();
if total_workers == 0 {
@@ -238,8 +374,11 @@ async fn run_benchmark(
let mut handles = Vec::new();
let total_hashes = Arc::new(std::sync::Mutex::new(0u64));
let cpu_chunk = 10_000u64;
let gpu_chunk = 1_000_000u64;
// Floor range widths at the old constants: engines still batch at the flag's
// size internally, but tiny flags don't turn per-call harness overhead into
// the measured quantity.
let cpu_chunk = cpu_batch_size.max(10_000);
let gpu_chunk = (gpu_batch_size as u64).max(1_000_000);
for worker_id in 0..total_workers {
let (engine, nonces_per_batch) = if worker_id < effective_cpu_workers {
@@ -255,31 +394,40 @@ async fn run_benchmark(
let handle = thread::spawn(move || {
let stride = U512::from(1_000_000_000_000u64);
let worker_start = U512::from(worker_id as u64).saturating_mul(stride);
let worker_range = EngineRange {
start: worker_start,
end: worker_start
.saturating_add(U512::from(nonces_per_batch))
.saturating_sub(U512::from(1u64)),
};
let mut current = U512::from(worker_id as u64).saturating_mul(stride);
let step = U512::from(nonces_per_batch);
loop {
if cancel.load(std::sync::atomic::Ordering::Relaxed) {
break;
}
let worker_range = EngineRange {
start: current,
end: current
.saturating_add(step)
.saturating_sub(U512::from(1u64)),
};
let cancel_check = AtomicBoolCancelCheck(&cancel);
let result = engine.search_range(&ctx, worker_range.clone(), &cancel_check);
let result = engine.search_range(&ctx, worker_range, &cancel_check);
match result {
engine_cpu::EngineStatus::Found { hash_count, .. }
| engine_cpu::EngineStatus::Exhausted { hash_count }
| engine_cpu::EngineStatus::Cancelled { hash_count } => {
| engine_cpu::EngineStatus::Cancelled { hash_count }
| engine_cpu::EngineStatus::DeviceLost { hash_count } => {
*hashes.lock().unwrap() += hash_count;
}
engine_cpu::EngineStatus::Running { .. } => {}
}
if matches!(result, engine_cpu::EngineStatus::DeviceLost { .. }) {
break;
}
current = current.saturating_add(step);
if start.elapsed() >= Duration::from_secs(duration) {
break;
}
crates/miner-service/Cargo.toml
+1 -1
View File
@@ -21,7 +21,7 @@ anyhow = { workspace = true }
# QUIC transport
quinn = "0.10"
rustls = { version = "0.21", default-features = false, features = ["dangerous_configuration", "quic"] }
quic-transport = { path = "../quic-transport" }
getrandom = "0.2"
# Protocol types from the node repo
crates/miner-service/src/lib.rs
+52 -3
View File
@@ -23,6 +23,10 @@ use std::thread;
pub struct ServiceConfig {
/// Address of the node to connect to (e.g., "127.0.0.1:9833").
pub node_addr: std::net::SocketAddr,
/// Shared secret that must match the node's miner auth token.
pub auth_token: String,
/// SHA-256 fingerprint (hex) of the node's miner TLS certificate DER.
pub tls_cert_sha256: String,
/// Number of CPU worker threads to use for mining (None = auto-detect)
pub cpu_workers: Option<usize>,
/// Number of GPU devices to use for mining (None = auto-detect)
@@ -33,6 +37,8 @@ pub struct ServiceConfig {
pub cpu_batch_size: u64,
/// GPU throttle delay in milliseconds between batches (0 = no throttle)
pub gpu_throttle_ms: u64,
/// Allow integrated GPUs even when discrete GPUs are available
pub allow_integrated: bool,
}
/// Engine type for tracking metrics per compute type.
@@ -206,12 +212,34 @@ impl WorkerPool {
// Dispatch job to all workers using bounded channels (capacity 16).
// Workers drain to get the latest job, so we just need room to queue.
let mut disconnected_count = 0;
for (i, tx) in self.job_senders.iter().enumerate() {
if let Err(e) = tx.try_send(job.clone()) {
match e {
crossbeam_channel::TrySendError::Disconnected(_) => {
// Worker thread has exited (e.g., device lost)
disconnected_count += 1;
log::debug!("Worker {i} channel disconnected (worker exited)");
}
crossbeam_channel::TrySendError::Full(_) => {
log::warn!(
"Failed to send job {new_job_id} to worker {i}: channel full - \
worker may be stuck or jobs arriving too fast"
);
}
}
}
}
if disconnected_count > 0 {
let active = self.job_senders.len() - disconnected_count;
if active == 0 {
log::error!(
"Failed to send job {new_job_id} to worker {i}: {e}. \
Channel full - worker may be stuck or jobs arriving too fast."
"All workers have exited! No workers available to process jobs. \
Consider restarting the miner."
);
} else {
log::warn!("{disconnected_count} worker(s) have exited, {active} still active");
}
}
@@ -336,6 +364,7 @@ fn worker_loop(
engine_cpu::EngineStatus::Found { .. } => "FOUND",
engine_cpu::EngineStatus::Exhausted { .. } => "EXHAUSTED",
engine_cpu::EngineStatus::Cancelled { .. } => "CANCELLED",
engine_cpu::EngineStatus::DeviceLost { .. } => "DEVICE_LOST",
engine_cpu::EngineStatus::Running { .. } => "RUNNING",
};
log::debug!(
@@ -352,6 +381,7 @@ fn worker_loop(
engine_cpu::EngineStatus::Found { hash_count, .. } => hash_count,
engine_cpu::EngineStatus::Exhausted { hash_count } => hash_count,
engine_cpu::EngineStatus::Cancelled { hash_count } => hash_count,
engine_cpu::EngineStatus::DeviceLost { hash_count } => hash_count,
engine_cpu::EngineStatus::Running { .. } => 0,
};
log_worker_completion(
@@ -400,6 +430,21 @@ fn worker_loop(
log_worker_completion(type_str, thread_id, "new block", hash_count, search_elapsed);
(None, hash_count)
}
engine_cpu::EngineStatus::DeviceLost { hash_count } => {
log::error!(
"{type_str} worker {thread_id} GPU device lost - worker exiting permanently"
);
// Send final result before exiting
let _ = result_tx.try_send(WorkerResult {
thread_id,
engine_type,
job_id,
candidate: None,
hash_count,
completed: true,
});
break; // Exit the worker loop
}
engine_cpu::EngineStatus::Running { .. } => {
// Should not happen for synchronous search
(None, 0)
@@ -430,6 +475,7 @@ pub fn resolve_gpu_configuration(
requested_devices: Option<usize>,
batch_size: u32,
throttle_ms: u64,
allow_integrated: bool,
) -> anyhow::Result<(Option<Arc<dyn MinerEngine>>, usize)> {
// Explicit 0 means no GPU
if requested_devices == Some(0) {
@@ -437,7 +483,7 @@ pub fn resolve_gpu_configuration(
}
// Try to initialize GPU engine
let engine = engine_gpu::GpuEngine::try_new(batch_size, throttle_ms);
let engine = engine_gpu::GpuEngine::try_new(batch_size, throttle_ms, allow_integrated);
let engine = match engine {
Ok(e) => e,
Err(e) => {
@@ -482,6 +528,7 @@ pub async fn run(config: ServiceConfig) -> anyhow::Result<()> {
config.gpu_devices,
config.gpu_batch_size,
config.gpu_throttle_ms,
config.allow_integrated,
)?;
// Resolve CPU workers
@@ -539,6 +586,8 @@ pub async fn run(config: ServiceConfig) -> anyhow::Result<()> {
log::info!("🌐 Connecting to node at {node_addr}");
quic::connect_and_mine(
config.node_addr,
&config.auth_token,
&config.tls_cert_sha256,
cpu_engine,
gpu_engine,
cpu_workers,
crates/miner-service/src/quic.rs
+50 -54
View File
@@ -10,8 +10,6 @@ use std::time::{Duration, Instant};
use engine_cpu::MinerEngine;
use primitive_types::U512;
use quinn::{ClientConfig, Endpoint};
use rustls::client::ServerCertVerified;
use quantus_miner_api::{
read_message, write_message, ApiResponseStatus, MinerMessage, MiningResult,
@@ -28,6 +26,8 @@ use pow_core::format_hashrate;
/// Uses a persistent worker pool to avoid thread creation overhead between jobs.
pub async fn connect_and_mine(
node_addr: SocketAddr,
auth_token: &str,
tls_cert_sha256: &str,
cpu_engine: Option<Arc<dyn MinerEngine>>,
gpu_engine: Option<Arc<dyn MinerEngine>>,
cpu_workers: usize,
@@ -42,18 +42,39 @@ pub async fn connect_and_mine(
loop {
log::info!("⛏️ Connecting to node at {}...", node_addr);
match establish_connection(node_addr).await {
match establish_connection(node_addr, auth_token, tls_cert_sha256).await {
Ok((connection, send, recv)) => {
log::info!("⛏️ Connected to node at {}", node_addr);
reconnect_delay = Duration::from_secs(1);
if let Err(e) = handle_connection(connection, send, recv, &worker_pool).await {
log::info!("⛏️ Connection lost: {}", e);
let mut authenticated = false;
if let Err(e) =
handle_connection(connection, send, recv, &worker_pool, &mut authenticated)
.await
{
// Cancel any running job when connection drops
worker_pool.cancel();
if e.downcast_ref::<quic_transport::PermanentConnectError>()
.is_some()
{
log::error!("⛏️ Permanent connection error (not retrying): {e}");
return Err(e);
}
log::info!("⛏️ Connection lost: {}", e);
}
// 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.
if authenticated {
reconnect_delay = Duration::from_secs(1);
}
}
Err(e) => {
if e.downcast_ref::<quic_transport::PermanentConnectError>()
.is_some()
{
log::error!("⛏️ Permanent connection error (not retrying): {e}");
return Err(e);
}
log::warn!("⛏️ Failed to connect to node: {}", e);
}
}
@@ -64,38 +85,18 @@ pub async fn connect_and_mine(
}
}
/// Establish a QUIC connection to the node.
/// Establish a QUIC connection to the node (shared transport crate).
async fn establish_connection(
addr: SocketAddr,
auth_token: &str,
tls_cert_sha256: &str,
) -> anyhow::Result<(quinn::Connection, quinn::SendStream, quinn::RecvStream)> {
let mut crypto = rustls::ClientConfig::builder()
.with_safe_defaults()
.with_custom_certificate_verifier(Arc::new(InsecureCertVerifier))
.with_no_client_auth();
crypto.alpn_protocols = vec![b"quantus-miner".to_vec()];
let mut client_config = ClientConfig::new(Arc::new(crypto));
let mut transport_config = quinn::TransportConfig::default();
transport_config.keep_alive_interval(Some(Duration::from_secs(5)));
transport_config.max_idle_timeout(Some(Duration::from_secs(15).try_into().unwrap()));
client_config.transport_config(Arc::new(transport_config));
let mut endpoint = Endpoint::client("0.0.0.0:0".parse().unwrap())?;
endpoint.set_default_client_config(client_config);
let connection = endpoint.connect(addr, "localhost")?.await?;
log::info!("⛏️ QUIC connection established to {}", addr);
log::info!("⛏️ Opening bidirectional stream to node...");
let (mut send, recv) = connection.open_bi().await?;
// Send Ready message to establish the stream
write_message(&mut send, &MinerMessage::Ready).await?;
log::info!("⛏️ Bidirectional stream established");
Ok((connection, send, recv))
let result = quic_transport::connect(addr, auth_token, tls_cert_sha256).await?;
log::info!(
"⛏️ QUIC connection and bidirectional stream established to {}",
addr
);
Ok(result)
}
/// Helper to send a message while monitoring connection health.
@@ -116,11 +117,15 @@ async fn send_message_checked(
}
/// Handle an established connection, receiving jobs and sending results.
///
/// Sets `authenticated` when the first `NewJob` arrives (proof the node
/// accepted our Ready token).
async fn handle_connection(
connection: quinn::Connection,
mut send: quinn::SendStream,
mut recv: quinn::RecvStream,
worker_pool: &WorkerPool,
authenticated: &mut bool,
) -> anyhow::Result<()> {
use crossbeam_channel::RecvTimeoutError;
@@ -220,12 +225,20 @@ async fn handle_connection(
biased;
reason = connection.closed() => {
return Err(anyhow::anyhow!("Connection closed: {}", reason));
let msg = reason.to_string();
if !*authenticated && msg.to_ascii_lowercase().contains("auth") {
return Err(quic_transport::PermanentConnectError(format!(
"node rejected miner auth ({msg}); check --auth-token / miner-auth-token"
))
.into());
}
return Err(anyhow::anyhow!("Connection closed: {}", msg));
}
msg_result = read_message(&mut recv) => {
match msg_result {
Ok(MinerMessage::NewJob(request)) => {
*authenticated = true;
log::info!(
"⛏️ Received job: id={}, hash=0x{}",
request.job_id,
@@ -305,7 +318,7 @@ async fn handle_connection(
Ok(MinerMessage::JobResult(_)) => {
log::warn!("Received unexpected JobResult from node");
}
Ok(MinerMessage::Ready) => {
Ok(MinerMessage::Ready { .. }) => {
log::warn!("Received unexpected Ready from node");
}
Err(e) if e.kind() == std::io::ErrorKind::UnexpectedEof => {
@@ -322,20 +335,3 @@ async fn handle_connection(
}
}
}
/// Certificate verifier that accepts any certificate (for self-signed certs).
struct InsecureCertVerifier;
impl rustls::client::ServerCertVerifier for InsecureCertVerifier {
fn verify_server_cert(
&self,
_end_entity: &rustls::Certificate,
_intermediates: &[rustls::Certificate],
_server_name: &rustls::ServerName,
_scts: &mut dyn Iterator<Item = &[u8]>,
_ocsp_response: &[u8],
_now: std::time::SystemTime,
) -> Result<ServerCertVerified, rustls::Error> {
Ok(ServerCertVerified::assertion())
}
}
crates/pool-service/Cargo.toml
+32
View File
@@ -0,0 +1,32 @@
[package]
name = "pool-service"
version.workspace = true
edition.workspace = true
publish = false
description = "Captcha share pool: turns browser proof-of-work captcha solves into Quantus mining shares"
[dependencies]
anyhow = { workspace = true }
clap = { workspace = true }
env_logger = { workspace = true }
hex = { workspace = true }
log = { workspace = true }
primitive-types = { workspace = true }
rand = { workspace = true }
serde = { workspace = true }
serde_json = { workspace = true, features = ["std"] }
tokio = { workspace = true }
warp = { workspace = true }
# QUIC transport (shared with miner-service)
quinn = "0.10"
quic-transport = { path = "../quic-transport" }
# Constant-time comparison for the site secret
constant_time_eq = "0.3"
# Protocol types from the node repo
quantus-miner-api = { workspace = true }
# Local crates
pow-core = { path = "../pow-core" }
crates/pool-service/src/http.rs
+219
View File
@@ -0,0 +1,219 @@
//! HTTP API for the captcha widget (session/share) and for protected sites
//! (siteverify), plus optional static file serving for the demo page.
use std::net::{IpAddr, SocketAddr};
use std::path::PathBuf;
use std::sync::Arc;
use primitive_types::U512;
use serde::{Deserialize, Serialize};
use warp::Filter;
use crate::rate_limit::SessionIssuerLimiter;
use crate::state::{PoolState, ShareOutcome};
#[derive(Serialize)]
struct SessionResponse {
session_id: String,
/// Hex, 32 bytes, no 0x prefix.
header_hash: String,
/// Hex U512, no 0x prefix. Solver iterates nonce_start..nonce_end.
nonce_start: String,
nonce_end: String,
/// Share target as hex U512: a share is valid iff hash < share_target.
/// (Target form avoids the solver needing U512 division.)
share_target: String,
/// Expected number of hashes to find a share (= share difficulty), FYI/UX.
expected_hashes: u64,
expires_in_secs: u64,
}
#[derive(Deserialize)]
struct ShareRequest {
session_id: String,
/// Hex U512 nonce, no 0x prefix.
nonce: String,
}
#[derive(Serialize)]
struct ShareResponse {
success: bool,
#[serde(skip_serializing_if = "Option::is_none")]
token: Option<String>,
#[serde(skip_serializing_if = "Option::is_none")]
error: Option<&'static str>,
block_found: bool,
}
/// Request shape mirrors reCAPTCHA/Turnstile: `secret` + `response`.
#[derive(Deserialize)]
struct SiteVerifyRequest {
secret: String,
response: String,
}
#[derive(Serialize)]
struct SiteVerifyResponse {
success: bool,
#[serde(rename = "challenge_ts", skip_serializing_if = "Option::is_none")]
challenge_ts: Option<u64>,
#[serde(rename = "error-codes", skip_serializing_if = "Vec::is_empty")]
error_codes: Vec<&'static str>,
}
pub async fn serve(
state: Arc<PoolState>,
limiter: Arc<SessionIssuerLimiter>,
addr: SocketAddr,
serve_dir: Option<PathBuf>,
) {
let with_state = {
let state = state.clone();
warp::any().map(move || state.clone())
};
let with_limiter = {
let limiter = limiter.clone();
warp::any().map(move || limiter.clone())
};
// POST /api/session -> new captcha challenge
let session = warp::path!("api" / "session")
.and(warp::post())
.and(warp::addr::remote())
.and(with_limiter.clone())
.and(with_state.clone())
.map(
|remote: Option<SocketAddr>,
limiter: Arc<SessionIssuerLimiter>,
state: Arc<PoolState>| {
let ip = remote
.map(|a| a.ip())
.unwrap_or(IpAddr::from([127, 0, 0, 1]));
if let Err(reason) = limiter.check(ip) {
return reply_error(warp::http::StatusCode::TOO_MANY_REQUESTS, reason);
}
match state.issue_session() {
Err(reason) => reply_error(warp::http::StatusCode::SERVICE_UNAVAILABLE, reason),
Ok(s) => {
let target = U512::MAX / s.share_difficulty;
let expected = s.share_difficulty.min(U512::from(u64::MAX)).as_u64();
warp::reply::with_status(
warp::reply::json(&SessionResponse {
session_id: s.id,
header_hash: hex::encode(s.job.header),
nonce_start: format!("{:x}", s.nonce_start),
nonce_end: format!("{:x}", s.nonce_end),
share_target: format!("{:x}", target),
expected_hashes: expected,
expires_in_secs: crate::state::SESSION_TTL.as_secs(),
}),
warp::http::StatusCode::OK,
)
}
}
},
);
// POST /api/share -> verify a solved nonce, mint a token
let share = warp::path!("api" / "share")
.and(warp::post())
.and(warp::body::content_length_limit(4096))
.and(warp::body::json())
.and(with_state.clone())
.map(|req: ShareRequest, state: Arc<PoolState>| {
let nonce = match U512::from_str_radix(req.nonce.trim_start_matches("0x"), 16) {
Ok(n) => n,
Err(_) => {
return warp::reply::with_status(
warp::reply::json(&ShareResponse {
success: false,
token: None,
error: Some("malformed_nonce"),
block_found: false,
}),
warp::http::StatusCode::BAD_REQUEST,
)
}
};
match state.submit_share(&req.session_id, nonce) {
ShareOutcome::Accepted { token, block_found } => warp::reply::with_status(
warp::reply::json(&ShareResponse {
success: true,
token: Some(token),
error: None,
block_found,
}),
warp::http::StatusCode::OK,
),
ShareOutcome::Rejected(reason) => warp::reply::with_status(
warp::reply::json(&ShareResponse {
success: false,
token: None,
error: Some(reason),
block_found: false,
}),
warp::http::StatusCode::FORBIDDEN,
),
}
});
// POST /siteverify -> protected site's backend redeems a token
let siteverify = warp::path!("siteverify")
.and(warp::post())
.and(warp::body::content_length_limit(4096))
// Accept both JSON and form encoding, like the incumbents do.
.and(
warp::body::json()
.or(warp::body::form())
.unify()
.map(|req: SiteVerifyRequest| req),
)
.and(with_state.clone())
.map(|req: SiteVerifyRequest, state: Arc<PoolState>| {
match state.verify_token(&req.secret, &req.response) {
Ok(ts) => warp::reply::json(&SiteVerifyResponse {
success: true,
challenge_ts: Some(ts),
error_codes: vec![],
}),
Err(code) => warp::reply::json(&SiteVerifyResponse {
success: false,
challenge_ts: None,
error_codes: vec![code],
}),
}
});
// GET /api/stats -> pool counters
let stats = warp::path!("api" / "stats")
.and(warp::get())
.and(with_state.clone())
.map(|state: Arc<PoolState>| warp::reply::json(&state.stats()));
let cors = warp::cors()
.allow_any_origin()
.allow_methods(vec!["GET", "POST", "OPTIONS"])
.allow_headers(vec!["content-type"]);
let api = session.or(share).or(siteverify).or(stats);
log::info!("HTTP API listening on {}", addr);
if let Some(dir) = serve_dir {
log::info!("Serving demo statics from {}", dir.display());
let statics = warp::get().and(warp::fs::dir(dir));
warp::serve(api.or(statics).with(cors)).run(addr).await;
} else {
warp::serve(api.with(cors)).run(addr).await;
}
}
fn reply_error(
status: warp::http::StatusCode,
error: &'static str,
) -> warp::reply::WithStatus<warp::reply::Json> {
#[derive(Serialize)]
struct Err<'a> {
error: &'a str,
}
warp::reply::with_status(warp::reply::json(&Err { error }), status)
}
crates/pool-service/src/main.rs
+250
View File
@@ -0,0 +1,250 @@
//! Quantus captcha share pool.
//!
//! Sits between a quantus-node (external miner protocol) and browser captcha
//! solvers: hands out low-difficulty share challenges over the real block
//! header, verifies solves, mints single-use tokens for site backends, and
//! submits any share that happens to meet full network difficulty as a block.
mod http;
mod rate_limit;
mod state;
mod upstream;
use std::net::SocketAddr;
use std::path::PathBuf;
use std::sync::Arc;
use std::time::Duration;
use clap::Parser;
use primitive_types::U512;
use crate::rate_limit::SessionRateLimit;
#[derive(Parser, Debug)]
#[command(name = "pool-service", about = "Quantus captcha share pool")]
struct Args {
/// HTTP listen address for the captcha/siteverify API.
#[arg(long, default_value = "127.0.0.1:8787", env = "POOL_HTTP_ADDR")]
http_addr: SocketAddr,
/// Address of the quantus-node external-miner QUIC endpoint.
/// If omitted, runs in standalone mode with synthetic jobs (demo only).
#[arg(long, env = "POOL_NODE_ADDR")]
node_addr: Option<SocketAddr>,
/// Shared auth token from the node's `miner-auth-token` file.
/// Required when `--node-addr` is set. Prefer `--auth-token-file`.
#[arg(long, env = "POOL_AUTH_TOKEN", conflicts_with = "auth_token_file")]
auth_token: Option<String>,
/// Path to the node's `miner-auth-token` file (trimmed).
#[arg(long, env = "POOL_AUTH_TOKEN_FILE", conflicts_with = "auth_token")]
auth_token_file: Option<PathBuf>,
/// SHA-256 fingerprint of the node's miner TLS certificate (64 hex chars).
/// Required when `--node-addr` is set. Prefer `--tls-cert-sha256-file`.
#[arg(
long,
env = "POOL_TLS_CERT_SHA256",
conflicts_with = "tls_cert_sha256_file"
)]
tls_cert_sha256: Option<String>,
/// Path to the node's `miner-tls-cert-sha256` file.
#[arg(
long,
env = "POOL_TLS_CERT_SHA256_FILE",
conflicts_with = "tls_cert_sha256"
)]
tls_cert_sha256_file: Option<PathBuf>,
/// Share difficulty: expected number of hashes per captcha solve.
/// Measured browser WASM rate ≈ 120 kH/s on an M-series laptop, so
/// 50000 ≈ 0.4 s desktop / ~2 s phone. Raise for stronger rate limiting.
#[arg(long, default_value = "50000", env = "POOL_SHARE_DIFFICULTY")]
share_difficulty: u64,
/// Secret that protected-site backends must present to /siteverify.
#[arg(long, default_value = "dev-secret", env = "POOL_SITE_SECRET")]
site_secret: String,
/// Optional directory of static files to serve (demo page / widget).
#[arg(long, env = "POOL_SERVE_DIR")]
serve_dir: Option<PathBuf>,
/// Job rotation interval in standalone mode, seconds.
#[arg(long, default_value = "20")]
standalone_job_secs: u64,
/// Maximum live captcha sessions; /api/session returns 503 when full.
#[arg(long, default_value = "100000", env = "POOL_MAX_SESSIONS")]
max_sessions: usize,
/// Maximum live (unredeemed) share tokens; shares are refused when full.
#[arg(long, default_value = "100000", env = "POOL_MAX_TOKENS")]
max_tokens: usize,
/// Max /api/session issuances per client IP per minute.
#[arg(long, default_value = "60", env = "POOL_SESSIONS_PER_IP_PER_MIN")]
sessions_per_ip_per_min: u32,
}
#[tokio::main]
async fn main() -> anyhow::Result<()> {
env_logger::Builder::from_env(env_logger::Env::default().default_filter_or("info")).init();
let args = Args::parse();
anyhow::ensure!(args.share_difficulty > 0, "--share-difficulty must be > 0");
if args.site_secret == "dev-secret" {
log::warn!("Using default site secret; set --site-secret in production");
}
// Resolve + validate node auth before starting HTTP / upstream tasks so a
// bad fingerprint or oversized token exits instead of retrying forever.
let node_upstream = match args.node_addr {
Some(addr) => {
let auth_token = resolve_auth_token(args.auth_token, args.auth_token_file)?;
let tls_cert_sha256 =
resolve_tls_cert_sha256(args.tls_cert_sha256, args.tls_cert_sha256_file)?;
quic_transport::validate_auth_config(&auth_token, &tls_cert_sha256)?;
Some((addr, auth_token, tls_cert_sha256))
}
None => {
if args.auth_token.is_some()
|| args.auth_token_file.is_some()
|| args.tls_cert_sha256.is_some()
|| args.tls_cert_sha256_file.is_some()
{
log::warn!("Ignoring auth/TLS pin flags in standalone mode (no --node-addr)");
}
None
}
};
let (solution_tx, solution_rx) = tokio::sync::mpsc::channel(16);
let state = state::PoolState::new(
U512::from(args.share_difficulty),
args.site_secret.clone(),
solution_tx,
state::Limits {
max_sessions: args.max_sessions,
max_tokens: args.max_tokens,
},
);
let limiter = Arc::new(rate_limit::SessionIssuerLimiter::new(SessionRateLimit {
max_per_ip: args.sessions_per_ip_per_min,
window: Duration::from_secs(60),
}));
// Periodic cleanup of expired sessions/tokens and stale rate-limit windows.
{
let state = state.clone();
let limiter = limiter.clone();
tokio::spawn(async move {
loop {
tokio::time::sleep(Duration::from_secs(30)).await;
state.gc();
limiter.gc();
}
});
}
// Job source. Node mode reconnects across transient disconnects (node
// restart); only a PermanentConnectError ends the task — then we tear
// down HTTP so the process does not look healthy on a bad token/pin.
match node_upstream {
Some((addr, auth_token, tls_cert_sha256)) => {
log::info!("Upstream: quantus-node at {}", addr);
let state_for_upstream = state.clone();
let upstream = tokio::spawn(async move {
upstream::run_node_client(
state_for_upstream,
addr,
auth_token,
tls_cert_sha256,
solution_rx,
)
.await
});
tokio::select! {
biased;
result = upstream => {
match result {
Ok(Ok(())) => anyhow::bail!(
"upstream task exited unexpectedly (should reconnect forever)"
),
Ok(Err(e)) => {
log::error!("Upstream permanent failure; shutting down: {e}");
Err(e.into())
}
Err(e) => Err(anyhow::anyhow!("upstream task panicked: {e}")),
}
}
_ = http::serve(state, limiter, args.http_addr, args.serve_dir) => Ok(()),
}
}
None => {
log::warn!("No --node-addr given: running STANDALONE with synthetic jobs");
let state_for_upstream = state.clone();
tokio::spawn(async move {
upstream::run_standalone(
state_for_upstream,
Duration::from_secs(args.standalone_job_secs),
solution_rx,
)
.await
});
http::serve(state, limiter, args.http_addr, args.serve_dir).await;
Ok(())
}
}
}
fn resolve_auth_token(
auth_token: Option<String>,
auth_token_file: Option<PathBuf>,
) -> anyhow::Result<String> {
resolve_required_secret(
auth_token,
auth_token_file,
"--auth-token",
"--auth-token-file",
"when --node-addr is set, pass --auth-token-file <PATH> to the node's \
miner-auth-token file (or --auth-token <TOKEN>)",
)
}
fn resolve_tls_cert_sha256(value: Option<String>, file: Option<PathBuf>) -> anyhow::Result<String> {
resolve_required_secret(
value,
file,
"--tls-cert-sha256",
"--tls-cert-sha256-file",
"when --node-addr is set, pass --tls-cert-sha256-file <PATH> to the node's \
miner-tls-cert-sha256 file (or --tls-cert-sha256 <HEX>; also printed in node logs)",
)
}
fn resolve_required_secret(
value: Option<String>,
file: Option<PathBuf>,
value_flag: &str,
file_flag: &str,
missing_msg: &str,
) -> anyhow::Result<String> {
if let Some(value) = value {
let value = value.trim().to_string();
anyhow::ensure!(!value.is_empty(), "{value_flag} is empty");
return Ok(value);
}
if let Some(path) = file {
let contents = std::fs::read_to_string(&path)
.map_err(|e| anyhow::anyhow!("failed to read {file_flag} {}: {}", path.display(), e))?;
let value = contents.trim().to_string();
anyhow::ensure!(!value.is_empty(), "{file_flag} {} is empty", path.display());
return Ok(value);
}
anyhow::bail!("{missing_msg}");
}
crates/pool-service/src/rate_limit.rs
+104
View File
@@ -0,0 +1,104 @@
//! Per-client rate limiting for unauthenticated endpoints.
//!
//! `/api/session` is free, so a hard map cap alone still lets an attacker
//! churn entries at the TTL boundary. A per-IP issuance ceiling bounds that
//! even when the global cap has headroom.
use std::collections::HashMap;
use std::net::IpAddr;
use std::sync::Mutex;
use std::time::{Duration, Instant};
#[derive(Debug, Clone, Copy)]
pub struct SessionRateLimit {
/// Max session issuances per IP per window.
pub max_per_ip: u32,
pub window: Duration,
}
impl Default for SessionRateLimit {
fn default() -> Self {
SessionRateLimit {
// 60 sessions/min/IP is generous for real users (a page load is
// one) but stops a single host from dominating issuance.
max_per_ip: 60,
window: Duration::from_secs(60),
}
}
}
#[derive(Default)]
struct IpWindow {
count: u32,
window_start: Option<Instant>,
}
pub struct SessionIssuerLimiter {
cfg: SessionRateLimit,
by_ip: Mutex<HashMap<IpAddr, IpWindow>>,
}
impl SessionIssuerLimiter {
pub fn new(cfg: SessionRateLimit) -> Self {
SessionIssuerLimiter {
cfg,
by_ip: Mutex::new(HashMap::new()),
}
}
/// Returns Ok(()) if this IP may issue another session, Err otherwise.
pub fn check(&self, ip: IpAddr) -> Result<(), &'static str> {
let now = Instant::now();
let mut map = self.by_ip.lock().unwrap();
let entry = map.entry(ip).or_default();
let window_start = entry.window_start.get_or_insert_with(Instant::now);
if now.duration_since(*window_start) >= self.cfg.window {
*window_start = now;
entry.count = 0;
}
if entry.count >= self.cfg.max_per_ip {
return Err("rate_limited");
}
entry.count += 1;
Ok(())
}
/// Drop stale IP windows so the side table doesn't grow forever.
pub fn gc(&self) {
let cutoff = self.cfg.window * 2;
let now = Instant::now();
self.by_ip.lock().unwrap().retain(|_, w| {
w.window_start
.map(|t| now.duration_since(t) < cutoff)
.unwrap_or(true)
});
}
}
#[cfg(test)]
mod tests {
use super::*;
use std::net::{Ipv4Addr, Ipv6Addr};
#[test]
fn allows_up_to_cap_then_rejects() {
let lim = SessionIssuerLimiter::new(SessionRateLimit {
max_per_ip: 2,
window: Duration::from_secs(60),
});
let ip = IpAddr::V4(Ipv4Addr::new(10, 0, 0, 1));
assert!(lim.check(ip).is_ok());
assert!(lim.check(ip).is_ok());
assert_eq!(lim.check(ip), Err("rate_limited"));
}
#[test]
fn ips_are_independent() {
let lim = SessionIssuerLimiter::new(SessionRateLimit {
max_per_ip: 1,
window: Duration::from_secs(60),
});
assert!(lim.check(IpAddr::V4(Ipv4Addr::new(1, 2, 3, 4))).is_ok());
assert!(lim.check(IpAddr::V6(Ipv6Addr::LOCALHOST)).is_ok());
}
}
crates/pool-service/src/state.rs
+527
View File
@@ -0,0 +1,527 @@
//! Core pool state: the current mining job, captcha sessions with disjoint
//! nonce ranges, and single-use share tokens.
use std::collections::HashMap;
use std::sync::atomic::{AtomicU64, Ordering};
use std::sync::{Arc, Mutex};
use std::time::{Duration, Instant, SystemTime, UNIX_EPOCH};
use primitive_types::U512;
use rand::RngCore;
/// How long a captcha session stays solvable after issuance.
pub const SESSION_TTL: Duration = Duration::from_secs(120);
/// How long a share token stays redeemable via /siteverify.
pub const TOKEN_TTL: Duration = Duration::from_secs(300);
/// Each session gets a disjoint slice of the nonce space this many nonces wide.
/// 2^64 nonces is unreachable within a session TTL at browser hash rates.
pub const SESSION_RANGE_BITS: u32 = 64;
/// Caps on live entries in the in-memory maps. `/api/session` is free and
/// unauthenticated, so without a ceiling a request flood becomes a memory-DoS
/// on the anti-abuse service itself. At the ~300 bytes/entry ballpark the
/// defaults bound each map to tens of MB.
#[derive(Debug, Clone, Copy)]
pub struct Limits {
pub max_sessions: usize,
pub max_tokens: usize,
}
impl Default for Limits {
fn default() -> Self {
Limits {
max_sessions: 100_000,
max_tokens: 100_000,
}
}
}
/// The job the pool is currently handing out to captcha sessions.
#[derive(Debug, Clone)]
pub struct Job {
/// Job id assigned by the upstream node (or synthesized in standalone mode).
pub job_id: String,
/// Block header hash to mine over.
pub header: [u8; 32],
/// Full network difficulty for this job (block found if a share meets it).
pub network_difficulty: U512,
}
/// A captcha session: one issued challenge, bound to a disjoint nonce range.
#[derive(Debug, Clone)]
pub struct Session {
pub id: String,
/// Snapshot of the job at issuance time. Shares are verified against this
/// header even if the pool has since moved to a newer job (grace period =
/// session TTL); only network-difficulty solutions for the *current* job
/// are submitted upstream.
pub job: Job,
pub nonce_start: U512,
pub nonce_end: U512, // exclusive
pub share_difficulty: U512,
pub issued_at: Instant,
pub solved: bool,
}
/// A verified share, redeemable exactly once by the protected site's backend.
#[derive(Debug, Clone)]
pub struct ShareToken {
pub created_at: Instant,
pub created_unix: u64,
pub consumed: bool,
}
#[derive(Debug, Default, Clone, serde::Serialize)]
pub struct PoolStats {
pub sessions_issued: u64,
pub shares_accepted: u64,
pub shares_rejected: u64,
pub tokens_verified: u64,
pub blocks_found: u64,
}
pub struct PoolState {
/// Current job, None until the first job arrives from the job source.
current_job: Mutex<Option<Job>>,
sessions: Mutex<HashMap<String, Session>>,
tokens: Mutex<HashMap<String, ShareToken>>,
stats: Mutex<PoolStats>,
/// Monotonic counter used to carve disjoint nonce ranges out of U512 space.
range_counter: AtomicU64,
/// Difficulty for captcha shares (expected hashes per solve).
pub share_difficulty: U512,
/// Site secret checked by /siteverify.
///
/// NOTE: single-tenant by design for now. All tokens are interchangeable
/// across any backend holding this secret; per-sitekey secrets land with
/// the host-registration follow-up.
pub site_secret: String,
limits: Limits,
/// Sender used to push network-difficulty solutions upstream.
pub solution_tx: tokio::sync::mpsc::Sender<FoundBlock>,
}
/// A share that met full network difficulty: a real block solution.
#[derive(Debug, Clone)]
pub struct FoundBlock {
pub job_id: String,
pub nonce: U512,
}
pub enum ShareOutcome {
/// Share accepted; token the client can hand to the protected site.
Accepted {
token: String,
block_found: bool,
},
Rejected(&'static str),
}
fn random_id() -> String {
let mut bytes = [0u8; 24];
rand::rng().fill_bytes(&mut bytes);
hex::encode(bytes)
}
fn unix_now() -> u64 {
SystemTime::now()
.duration_since(UNIX_EPOCH)
.map(|d| d.as_secs())
.unwrap_or(0)
}
impl PoolState {
pub fn new(
share_difficulty: U512,
site_secret: String,
solution_tx: tokio::sync::mpsc::Sender<FoundBlock>,
limits: Limits,
) -> Arc<Self> {
assert!(
!share_difficulty.is_zero(),
"share difficulty must be non-zero"
);
Arc::new(PoolState {
current_job: Mutex::new(None),
sessions: Mutex::new(HashMap::new()),
tokens: Mutex::new(HashMap::new()),
stats: Mutex::new(PoolStats::default()),
range_counter: AtomicU64::new(0),
share_difficulty,
site_secret,
limits,
solution_tx,
})
}
/// Install a new job (from the upstream node or the standalone generator).
/// Existing sessions keep their job snapshot and remain solvable until TTL.
pub fn set_job(&self, job: Job) {
log::info!(
"New job {}: header=0x{} network_difficulty={}",
job.job_id,
hex::encode(job.header),
pow_core::format_u512(job.network_difficulty)
);
*self.current_job.lock().unwrap() = Some(job);
}
pub fn current_job(&self) -> Option<Job> {
self.current_job.lock().unwrap().clone()
}
/// Drop the current job (e.g. upstream disconnected). New sessions will
/// see `no_job_available` until the next `set_job`.
pub fn clear_job(&self) {
if self.current_job.lock().unwrap().take().is_some() {
log::info!("Cleared current job (upstream unavailable)");
}
}
/// Issue a new captcha session with a disjoint nonce range.
pub fn issue_session(&self) -> Result<Session, &'static str> {
let job = self.current_job().ok_or("no_job_available")?;
if self.sessions.lock().unwrap().len() >= self.limits.max_sessions {
return Err("at_capacity");
}
// Carve the next disjoint 2^64-nonce slice. The counter is 64 bits and
// slices are 2^64 wide, so slices live in the bottom 2^128 of the U512
// nonce space and never collide with each other.
let index = self.range_counter.fetch_add(1, Ordering::Relaxed);
let nonce_start = U512::from(index) << SESSION_RANGE_BITS;
let nonce_end = nonce_start + (U512::one() << SESSION_RANGE_BITS);
let session = Session {
id: random_id(),
job,
nonce_start,
nonce_end,
share_difficulty: self.share_difficulty,
issued_at: Instant::now(),
solved: false,
};
self.sessions
.lock()
.unwrap()
.insert(session.id.clone(), session.clone());
self.stats.lock().unwrap().sessions_issued += 1;
Ok(session)
}
/// Verify a submitted share and, if valid, mint a single-use token.
pub fn submit_share(&self, session_id: &str, nonce: U512) -> ShareOutcome {
let session = {
let mut sessions = self.sessions.lock().unwrap();
match sessions.get_mut(session_id) {
None => return self.reject("unknown_session"),
Some(s) => {
if s.solved {
return self.reject("session_already_solved");
}
if s.issued_at.elapsed() > SESSION_TTL {
sessions.remove(session_id);
return self.reject("session_expired");
}
// Mark solved optimistically; unmark on verification failure
// would allow retries, but a failed share means a buggy or
// malicious client, so we burn the session either way.
s.solved = true;
s.clone()
}
}
};
if nonce < session.nonce_start || nonce >= session.nonce_end {
return self.reject("nonce_out_of_range");
}
let (share_valid, hash) = pow_core::is_valid_nonce(
session.job.header,
nonce.to_big_endian(),
session.share_difficulty,
);
if !share_valid {
return self.reject("share_below_target");
}
// Jackpot check: does this share meet full network difficulty for the
// job it was issued against? Only submit if that job is still current.
// Target derivation delegates to pow-core's canonical JobContext.
let network_target =
pow_core::JobContext::new(session.job.header, session.job.network_difficulty).target;
let mut block_found = false;
if hash < network_target {
let still_current = self
.current_job()
.map(|j| j.job_id == session.job.job_id)
.unwrap_or(false);
if still_current {
log::info!(
"BLOCK FOUND by captcha share! job={} nonce={:x}",
session.job.job_id,
nonce
);
// This is the revenue event: count it and be loud if the
// upstream queue can't take it.
match self.solution_tx.try_send(FoundBlock {
job_id: session.job.job_id.clone(),
nonce,
}) {
Ok(()) => {
block_found = true;
self.stats.lock().unwrap().blocks_found += 1;
}
Err(e) => {
log::error!(
"DROPPED block solution for job {}: upstream queue unavailable ({})",
session.job.job_id,
e
);
}
}
}
}
let token = random_id();
{
let mut tokens = self.tokens.lock().unwrap();
// A full token map means someone is farming shares faster than
// TTL expiry; refuse rather than grow without bound. The client
// paid hashes for this, so make the refusal explicit.
if tokens.len() >= self.limits.max_tokens {
return self.reject("at_capacity");
}
tokens.insert(
token.clone(),
ShareToken {
created_at: Instant::now(),
created_unix: unix_now(),
consumed: false,
},
);
}
self.stats.lock().unwrap().shares_accepted += 1;
ShareOutcome::Accepted { token, block_found }
}
/// Redeem a share token (the reCAPTCHA/Turnstile "siteverify" semantics):
/// valid exactly once, only with the correct site secret.
pub fn verify_token(&self, secret: &str, token: &str) -> Result<u64, &'static str> {
if !constant_time_eq::constant_time_eq(secret.as_bytes(), self.site_secret.as_bytes()) {
return Err("invalid_secret");
}
let mut tokens = self.tokens.lock().unwrap();
match tokens.get_mut(token) {
None => Err("invalid_token"),
Some(t) if t.consumed => Err("token_already_consumed"),
Some(t) if t.created_at.elapsed() > TOKEN_TTL => {
tokens.remove(token);
Err("token_expired")
}
Some(t) => {
t.consumed = true;
let ts = t.created_unix;
self.stats.lock().unwrap().tokens_verified += 1;
Ok(ts)
}
}
}
pub fn stats(&self) -> PoolStats {
self.stats.lock().unwrap().clone()
}
/// Drop expired sessions and tokens. Called periodically.
pub fn gc(&self) {
self.sessions
.lock()
.unwrap()
.retain(|_, s| s.issued_at.elapsed() <= SESSION_TTL);
self.tokens
.lock()
.unwrap()
.retain(|_, t| t.created_at.elapsed() <= TOKEN_TTL);
}
fn reject(&self, reason: &'static str) -> ShareOutcome {
self.stats.lock().unwrap().shares_rejected += 1;
ShareOutcome::Rejected(reason)
}
}
#[cfg(test)]
mod tests {
use super::*;
fn test_state(
share_difficulty: u64,
) -> (Arc<PoolState>, tokio::sync::mpsc::Receiver<FoundBlock>) {
let (tx, rx) = tokio::sync::mpsc::channel(4);
let state = PoolState::new(
U512::from(share_difficulty),
"secret".into(),
tx,
Limits::default(),
);
state.set_job(Job {
job_id: "1".into(),
header: [7u8; 32],
network_difficulty: U512::MAX, // effectively unreachable
});
(state, rx)
}
/// Brute-force a valid share within the session's range.
fn solve(session: &Session) -> U512 {
let mut nonce = session.nonce_start;
loop {
let (ok, _) = pow_core::is_valid_nonce(
session.job.header,
nonce.to_big_endian(),
session.share_difficulty,
);
if ok {
return nonce;
}
nonce += U512::one();
}
}
#[test]
fn accepts_valid_share_and_verifies_token_once() {
let (state, _rx) = test_state(2); // ~2 hashes per solve
let session = state.issue_session().unwrap();
let nonce = solve(&session);
let token = match state.submit_share(&session.id, nonce) {
ShareOutcome::Accepted { token, block_found } => {
assert!(!block_found);
token
}
ShareOutcome::Rejected(r) => panic!("rejected: {r}"),
};
assert!(state.verify_token("secret", &token).is_ok());
assert_eq!(
state.verify_token("secret", &token),
Err("token_already_consumed")
);
assert_eq!(state.verify_token("wrong", &token), Err("invalid_secret"));
}
#[test]
fn rejects_out_of_range_nonce() {
let (state, _rx) = test_state(1); // difficulty 1: every nonce is a valid hash
let session = state.issue_session().unwrap();
// A nonce outside the assigned slice must be rejected even though its
// hash trivially meets the share target.
let outside = session.nonce_end;
match state.submit_share(&session.id, outside) {
ShareOutcome::Rejected("nonce_out_of_range") => {}
other => panic!(
"expected out-of-range rejection, got {:?}",
matches_str(&other)
),
}
}
#[test]
fn burns_session_after_one_submission() {
let (state, _rx) = test_state(1);
let session = state.issue_session().unwrap();
let nonce = session.nonce_start;
assert!(matches!(
state.submit_share(&session.id, nonce),
ShareOutcome::Accepted { .. }
));
match state.submit_share(&session.id, nonce) {
ShareOutcome::Rejected("session_already_solved") => {}
other => panic!("expected burned session, got {:?}", matches_str(&other)),
}
}
#[test]
fn disjoint_ranges() {
let (state, _rx) = test_state(1);
let a = state.issue_session().unwrap();
let b = state.issue_session().unwrap();
assert!(a.nonce_end <= b.nonce_start || b.nonce_end <= a.nonce_start);
}
#[test]
fn block_found_signalled_when_share_meets_network_difficulty() {
let (tx, mut rx) = tokio::sync::mpsc::channel(4);
// Network difficulty == share difficulty == 2: any valid share is a block.
let state = PoolState::new(U512::from(2u64), "secret".into(), tx, Limits::default());
state.set_job(Job {
job_id: "1".into(),
header: [9u8; 32],
network_difficulty: U512::from(2u64),
});
let session = state.issue_session().unwrap();
let nonce = solve(&session);
match state.submit_share(&session.id, nonce) {
ShareOutcome::Accepted { block_found, .. } => assert!(block_found),
ShareOutcome::Rejected(r) => panic!("rejected: {r}"),
}
let block = rx
.try_recv()
.expect("block solution should be queued upstream");
assert_eq!(block.job_id, "1");
assert_eq!(block.nonce, nonce);
}
#[test]
fn session_cap_rejects_flood() {
let (tx, _rx) = tokio::sync::mpsc::channel(4);
let state = PoolState::new(
U512::from(1u64),
"secret".into(),
tx,
Limits {
max_sessions: 2,
max_tokens: 2,
},
);
state.set_job(Job {
job_id: "1".into(),
header: [7u8; 32],
network_difficulty: U512::MAX,
});
assert!(state.issue_session().is_ok());
assert!(state.issue_session().is_ok());
assert!(matches!(state.issue_session(), Err("at_capacity")));
}
#[test]
fn dropped_block_does_not_count_or_flag() {
// Zero-capacity channel: try_send always fails, simulating a stalled
// upstream. The share must still be accepted (the user did the work),
// but block_found must be false and blocks_found must stay 0.
let (tx, rx) = tokio::sync::mpsc::channel(1);
drop(rx); // closed channel -> try_send errors
let state = PoolState::new(U512::from(2u64), "secret".into(), tx, Limits::default());
state.set_job(Job {
job_id: "1".into(),
header: [9u8; 32],
network_difficulty: U512::from(2u64),
});
let session = state.issue_session().unwrap();
let nonce = solve(&session);
match state.submit_share(&session.id, nonce) {
ShareOutcome::Accepted { block_found, .. } => assert!(!block_found),
ShareOutcome::Rejected(r) => panic!("rejected: {r}"),
}
assert_eq!(state.stats().blocks_found, 0);
}
fn matches_str(o: &ShareOutcome) -> &'static str {
match o {
ShareOutcome::Accepted { .. } => "accepted",
ShareOutcome::Rejected(r) => r,
}
}
}
crates/pool-service/src/upstream.rs
+192
View File
@@ -0,0 +1,192 @@
//! Job sources: either a QUIC connection to a real quantus-node (speaking the
//! external miner protocol), or a standalone generator for demos/tests.
use std::net::SocketAddr;
use std::sync::Arc;
use std::time::Duration;
use primitive_types::U512;
use quantus_miner_api::{
read_message, write_message, ApiResponseStatus, MinerMessage, MiningResult,
};
use tokio::sync::mpsc::Receiver;
use crate::state::{FoundBlock, Job, PoolState};
/// Connect to a node as an external miner and keep the pool's job current.
/// Solutions arriving on `solutions` are pushed upstream as JobResults.
///
/// Transient disconnects (node restart, network blip) clear the current job
/// and reconnect with backoff — the pool process stays up. Only
/// [`quic_transport::PermanentConnectError`] (bad pin / rejected auth) returns
/// `Err` so `main` can shut down the HTTP side.
pub async fn run_node_client(
state: Arc<PoolState>,
node_addr: SocketAddr,
auth_token: String,
tls_cert_sha256: String,
mut solutions: Receiver<FoundBlock>,
) -> Result<(), quic_transport::PermanentConnectError> {
let mut reconnect_delay = Duration::from_secs(1);
const MAX_RECONNECT_DELAY: Duration = Duration::from_secs(30);
// A block taken off the channel but not yet acknowledged by a successful
// upstream write. Kept across reconnects so a failed write doesn't drop
// the block (the node may still accept it after reconnection if the job
// hasn't moved on).
let mut pending: Option<FoundBlock> = None;
loop {
log::info!("Connecting to node at {}...", node_addr);
match quic_transport::connect(node_addr, &auth_token, &tls_cert_sha256).await {
Ok((connection, mut send, mut recv)) => {
log::info!("Connected to node at {}", node_addr);
let mut authenticated = false;
loop {
// Retry any unsubmitted block before doing anything else.
if let Some(block) = pending.take() {
match submit_block(&mut send, &block).await {
Ok(()) => {
log::info!(
"Submitted block solution to node (job {})",
block.job_id
);
}
Err(e) => {
log::error!("Failed to submit block solution, will retry: {}", e);
pending = Some(block);
break; // reconnect and retry
}
}
}
tokio::select! {
biased;
reason = connection.closed() => {
let msg = reason.to_string();
if !authenticated && msg.to_ascii_lowercase().contains("auth") {
state.clear_job();
return Err(quic_transport::PermanentConnectError(format!(
"node rejected miner auth ({msg}); check --auth-token / miner-auth-token"
)));
}
log::warn!("Node connection closed: {}", msg);
break;
}
// A captcha share met network difficulty: submit it.
block = solutions.recv() => {
if let Some(block) = block {
// Park it in `pending`; the top of the loop
// submits it and keeps it on write failure.
pending = Some(block);
}
}
msg = read_message(&mut recv) => {
match msg {
Ok(MinerMessage::NewJob(request)) => {
authenticated = true;
match parse_job(&request.job_id, &request.mining_hash, &request.difficulty) {
Ok(job) => state.set_job(job),
Err(e) => log::warn!("Ignoring malformed job from node: {}", e),
}
}
Ok(other) => {
log::warn!("Unexpected message from node: {:?}", other);
}
Err(e) => {
log::warn!("Read error from node: {}", e);
break;
}
}
}
}
}
// Node gone / stream dropped: stop handing out captchas on a
// stale header until we reconnect and get a fresh NewJob.
state.clear_job();
if authenticated {
reconnect_delay = Duration::from_secs(1);
}
}
Err(e) => {
if let Some(perm) = e.downcast_ref::<quic_transport::PermanentConnectError>() {
state.clear_job();
return Err(quic_transport::PermanentConnectError(perm.0.clone()));
}
log::warn!("Failed to connect to node: {}", e);
state.clear_job();
}
}
log::info!("Reconnecting in {:?}...", reconnect_delay);
tokio::time::sleep(reconnect_delay).await;
reconnect_delay = (reconnect_delay * 2).min(MAX_RECONNECT_DELAY);
}
}
async fn submit_block(send: &mut quinn::SendStream, block: &FoundBlock) -> std::io::Result<()> {
let result = MiningResult {
status: ApiResponseStatus::Completed,
job_id: block.job_id.clone(),
nonce: Some(format!("{:x}", block.nonce)),
work: Some(hex::encode(block.nonce.to_big_endian())),
hash_count: 0,
elapsed_time: 0.0,
miner_id: None,
};
write_message(send, &MinerMessage::JobResult(result)).await
}
fn parse_job(job_id: &str, mining_hash: &str, difficulty: &str) -> anyhow::Result<Job> {
let bytes = hex::decode(mining_hash)?;
let header: [u8; 32] = bytes
.try_into()
.map_err(|_| anyhow::anyhow!("mining_hash must be 32 bytes"))?;
let network_difficulty =
U512::from_dec_str(difficulty).map_err(|e| anyhow::anyhow!("bad difficulty: {:?}", e))?;
if network_difficulty.is_zero() {
anyhow::bail!("zero difficulty");
}
Ok(Job {
job_id: job_id.to_string(),
header,
network_difficulty,
})
}
/// Standalone mode: synthesize a fresh random job every `interval` so the
/// captcha works without a running node. Solutions are logged and dropped.
pub async fn run_standalone(
state: Arc<PoolState>,
interval: Duration,
mut solutions: Receiver<FoundBlock>,
) {
use rand::RngCore;
let mut counter: u64 = 0;
loop {
let mut header = [0u8; 32];
rand::rng().fill_bytes(&mut header);
counter += 1;
state.set_job(Job {
job_id: format!("standalone-{}", counter),
header,
// Unreachable "network" difficulty: standalone mode never finds blocks.
network_difficulty: U512::MAX,
});
tokio::select! {
_ = tokio::time::sleep(interval) => {}
block = solutions.recv() => {
if let Some(block) = block {
log::info!("(standalone) would submit block for job {}", block.job_id);
}
}
}
}
}
crates/pow-core/Cargo.toml
+1
View File
@@ -34,6 +34,7 @@ gpu-interop = []
[dependencies]
primitive-types = { workspace = true }
qp-poseidon-core = { workspace = true }
qpow-math = { workspace = true }
[dev-dependencies]
crates/pow-core/src/lib.rs
+54
View File
@@ -1,7 +1,28 @@
use primitive_types::U512;
use qp_poseidon_core::{Goldilocks, Poseidon2};
pub use qp_poseidon_core::SPONGE_WIDTH;
pub use qpow_math::{get_nonce_hash, is_valid_nonce, mine_range};
/// Sponge state (canonical u64 felts) after absorbing the 32-byte header and the
/// high 32 bytes of the big-endian nonce — the first two of the five Poseidon2
/// permutations of `get_nonce_hash`. This state is identical for every nonce in a
/// batch as long as incrementing the nonce never carries into its high 256 bits,
/// so GPU kernels can resume the sponge from here and skip 2 of 5 permutations.
pub fn mining_midstate(header: [u8; 32], nonce_high_be: [u8; 32]) -> [u64; SPONGE_WIDTH] {
let poseidon2 = Poseidon2::new();
let mut state = [Goldilocks::ZERO; SPONGE_WIDTH];
for (i, chunk) in header.chunks_exact(4).enumerate() {
state[i] += Goldilocks::from_u64(u32::from_le_bytes(chunk.try_into().unwrap()) as u64);
}
poseidon2.permute_mut(&mut state);
for (i, chunk) in nonce_high_be.chunks_exact(4).enumerate() {
state[i] += Goldilocks::from_u64(u32::from_le_bytes(chunk.try_into().unwrap()) as u64);
}
poseidon2.permute_mut(&mut state);
state.map(|g| g.as_canonical_u64())
}
/// Format a U512 in a human-readable way (scientific notation for large numbers).
pub fn format_u512(n: U512) -> String {
if n.is_zero() {
@@ -222,6 +243,39 @@ mod tests {
assert!(result.is_none());
}
#[test]
fn test_midstate_resumes_to_full_hash() {
let header = [7u8; 32];
let nonce = (U512::from(0xdeadbeefcafeu64) << 300) | U512::from(0x1234567890abcdefu64);
let nonce_be = nonce.to_big_endian();
let mut state =
mining_midstate(header, nonce_be[..32].try_into().unwrap()).map(Goldilocks::from_u64);
let poseidon2 = Poseidon2::new();
for (i, chunk) in nonce_be[32..].chunks_exact(4).enumerate() {
state[i] += Goldilocks::from_u64(u32::from_le_bytes(chunk.try_into().unwrap()) as u64);
}
poseidon2.permute_mut(&mut state);
state[0] += Goldilocks::ONE;
state[1] += Goldilocks::ONE;
poseidon2.permute_mut(&mut state);
let mut hash = [0u8; 64];
for i in 0..4 {
hash[i * 8..(i + 1) * 8].copy_from_slice(&state[i].as_canonical_u64().to_le_bytes());
}
poseidon2.permute_mut(&mut state);
for i in 0..4 {
hash[32 + i * 8..32 + (i + 1) * 8]
.copy_from_slice(&state[i].as_canonical_u64().to_le_bytes());
}
assert_eq!(
U512::from_big_endian(&hash),
qpow_math::get_nonce_hash(header, nonce_be)
);
}
#[test]
#[should_panic(expected = "division by zero")]
fn test_zero_difficulty_panics() {
crates/quic-transport/Cargo.toml
+15
View File
@@ -0,0 +1,15 @@
[package]
name = "quic-transport"
version.workspace = true
edition.workspace = true
publish = false
description = "Shared QUIC client transport for connecting to a quantus-node's external-miner endpoint"
[dependencies]
anyhow = { workspace = true }
quantus-miner-api = { workspace = true }
quinn = "0.10"
rustls = { version = "0.21", default-features = false, features = ["dangerous_configuration", "quic"] }
serde_json = { workspace = true }
sha2 = "0.10"
tokio = { workspace = true, features = ["macros", "time"] }
crates/quic-transport/src/lib.rs
+218
View File
@@ -0,0 +1,218 @@
//! Shared QUIC client used by miner-service and pool-service to connect to a
//! quantus-node's external-miner endpoint.
//!
//! The node uses a persisted self-signed certificate. Callers must supply the
//! expected SHA-256 fingerprint of that certificate (lowercase hex); the
//! transport rejects any other server cert.
use std::net::SocketAddr;
use std::sync::Arc;
use std::time::Duration;
use quantus_miner_api::{write_message, MinerMessage, MAX_MESSAGE_SIZE};
use quinn::{ClientConfig, Endpoint};
use rustls::client::ServerCertVerified;
use sha2::{Digest, Sha256};
/// Permanent misconfiguration / auth failure: callers must not reconnect-loop.
#[derive(Debug)]
pub struct PermanentConnectError(pub String);
impl std::fmt::Display for PermanentConnectError {
fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
write!(f, "{}", self.0)
}
}
impl std::error::Error for PermanentConnectError {}
/// Parse and validate miner auth inputs that never become valid by retrying.
///
/// Call this at process startup (before workers / HTTP) so a bad fingerprint
/// or oversized token exits immediately instead of looping forever.
pub fn validate_auth_config(auth_token: &str, tls_cert_sha256_hex: &str) -> anyhow::Result<()> {
parse_fingerprint(tls_cert_sha256_hex)?;
validate_ready_frame(auth_token)?;
Ok(())
}
/// Parse a 64-hex (optional whitespace/`:` separators) SHA-256 fingerprint.
pub fn parse_fingerprint(tls_cert_sha256_hex: &str) -> anyhow::Result<[u8; 32]> {
normalize_fingerprint(tls_cert_sha256_hex).ok_or_else(|| {
PermanentConnectError("invalid TLS cert SHA-256 fingerprint (expected 64 hex chars)".into())
.into()
})
}
fn validate_ready_frame(auth_token: &str) -> anyhow::Result<()> {
let frame = serde_json::to_vec(&MinerMessage::Ready {
token: auth_token.to_string(),
})
.map_err(|e| PermanentConnectError(format!("failed to serialize Ready frame: {e}")))?;
if frame.len() > MAX_MESSAGE_SIZE as usize {
return Err(PermanentConnectError(format!(
"auth token serializes to a {}-byte Ready frame; max is {} \
(shorten the token or avoid characters that require JSON escaping)",
frame.len(),
MAX_MESSAGE_SIZE
))
.into());
}
Ok(())
}
/// Establish a QUIC connection to the node's external-miner endpoint, open
/// the bidirectional stream, and send the initial `Ready { token }` message.
///
/// `tls_cert_sha256_hex` is the lowercase hex SHA-256 of the node's miner TLS
/// certificate DER (contents of the node's `miner-tls-cert-sha256` file).
///
/// Wrong tokens are rejected by the node closing with "auth failed"; that is
/// returned as [`PermanentConnectError`] so callers do not reconnect forever.
pub async fn connect(
addr: SocketAddr,
auth_token: &str,
tls_cert_sha256_hex: &str,
) -> anyhow::Result<(quinn::Connection, quinn::SendStream, quinn::RecvStream)> {
// Re-check here so library callers that skip startup validation still fail
// closed instead of looping on a deterministic parse error.
let expected = parse_fingerprint(tls_cert_sha256_hex)?;
validate_ready_frame(auth_token)?;
let mut crypto = rustls::ClientConfig::builder()
.with_safe_defaults()
.with_custom_certificate_verifier(Arc::new(PinnedCertVerifier { expected }))
.with_no_client_auth();
// Versioned with the wire protocol (node side: MINER_ALPN). `/2` = the
// authenticated `Ready { token }` protocol; a mismatched node/miner pair
// fails at the TLS handshake with "no application protocol" instead of an
// opaque auth/deserialize error.
crypto.alpn_protocols = vec![b"quantus-miner/2".to_vec()];
let mut client_config = ClientConfig::new(Arc::new(crypto));
let mut transport_config = quinn::TransportConfig::default();
transport_config.keep_alive_interval(Some(Duration::from_secs(5)));
transport_config.max_idle_timeout(Some(Duration::from_secs(15).try_into().unwrap()));
client_config.transport_config(Arc::new(transport_config));
let mut endpoint = Endpoint::client("0.0.0.0:0".parse().unwrap())?;
endpoint.set_default_client_config(client_config);
let connection = match endpoint.connect(addr, "localhost")?.await {
Ok(c) => c,
Err(e) => {
let msg = e.to_string();
// Wrong pin fails the TLS handshake deterministically.
if msg.to_ascii_lowercase().contains("fingerprint") {
return Err(PermanentConnectError(format!(
"TLS certificate pin rejected ({msg}); check --tls-cert-sha256 / \
miner-tls-cert-sha256"
))
.into());
}
return Err(e.into());
}
};
let (mut send, recv) = connection.open_bi().await?;
write_message(
&mut send,
&MinerMessage::Ready {
token: auth_token.to_string(),
},
)
.await?;
// Ready is fire-and-forget on the wire; the node closes immediately on a
// bad token. Wait briefly so that surfaces as a permanent error here
// instead of "connected" + 1s reconnect churn.
const AUTH_REJECT_GRACE: Duration = Duration::from_secs(2);
tokio::select! {
reason = connection.closed() => {
let msg = reason.to_string();
if msg.to_ascii_lowercase().contains("auth") {
return Err(PermanentConnectError(format!(
"node rejected miner auth ({msg}); check --auth-token / miner-auth-token"
))
.into());
}
return Err(anyhow::anyhow!("connection closed during auth handshake: {msg}"));
}
_ = tokio::time::sleep(AUTH_REJECT_GRACE) => {}
}
Ok((connection, send, recv))
}
fn normalize_fingerprint(input: &str) -> Option<[u8; 32]> {
let hex: String = input
.chars()
.filter(|c| !c.is_whitespace() && *c != ':')
.map(|c| c.to_ascii_lowercase())
.collect();
if hex.len() != 64 || !hex.chars().all(|c| c.is_ascii_hexdigit()) {
return None;
}
let mut out = [0u8; 32];
for i in 0..32 {
out[i] = u8::from_str_radix(&hex[i * 2..i * 2 + 2], 16).ok()?;
}
Some(out)
}
/// Accepts only a server certificate whose SHA-256 fingerprint matches.
struct PinnedCertVerifier {
expected: [u8; 32],
}
impl rustls::client::ServerCertVerifier for PinnedCertVerifier {
fn verify_server_cert(
&self,
end_entity: &rustls::Certificate,
_intermediates: &[rustls::Certificate],
_server_name: &rustls::ServerName,
_scts: &mut dyn Iterator<Item = &[u8]>,
_ocsp_response: &[u8],
_now: std::time::SystemTime,
) -> Result<ServerCertVerified, rustls::Error> {
let actual = Sha256::digest(&end_entity.0);
if actual.as_slice() != self.expected.as_slice() {
return Err(rustls::Error::General(
"server certificate SHA-256 fingerprint mismatch".into(),
));
}
Ok(ServerCertVerified::assertion())
}
}
#[cfg(test)]
mod tests {
use super::*;
#[test]
fn rejects_short_fingerprint() {
let err = validate_auth_config("token", "nope").unwrap_err();
assert!(
err.downcast_ref::<PermanentConnectError>().is_some(),
"{err}"
);
}
#[test]
fn accepts_valid_fingerprint_and_token() {
let fp = "a".repeat(64);
validate_auth_config("deadbeef", &fp).unwrap();
}
#[test]
fn rejects_token_whose_ready_frame_exceeds_limit() {
let fp = "a".repeat(64);
// Escape-heavy token: raw length under a naive cap can still blow the frame.
let token = "\"".repeat(600);
let err = validate_auth_config(&token, &fp).unwrap_err();
assert!(
err.downcast_ref::<PermanentConnectError>().is_some(),
"{err}"
);
assert!(err.to_string().contains("Ready frame"), "{err}");
}
}
crates/solver-wasm/Cargo.toml
+13
View File
@@ -0,0 +1,13 @@
[package]
name = "solver-wasm"
version.workspace = true
edition.workspace = true
publish = false
description = "Browser WASM solver for Quantus captcha shares (raw C ABI, no wasm-bindgen)"
[lib]
crate-type = ["cdylib", "rlib"]
[dependencies]
primitive-types = { workspace = true }
qpow-math = { workspace = true }
crates/solver-wasm/src/lib.rs
+101
View File
@@ -0,0 +1,101 @@
//! Browser captcha solver: grinds Poseidon2 nonce hashes over a block header.
//!
//! Exposed as a raw C-ABI WASM module (no wasm-bindgen) so it can be built
//! with a bare `cargo build --target wasm32-unknown-unknown` and loaded with
//! minimal JS glue. All I/O goes through one static buffer:
//!
//! ```text
//! offset 0..32 header hash (in)
//! offset 32..96 nonce, 64-byte big-endian (in: start; out: found/next)
//! offset 96..160 share target, 64-byte BE (in)
//! offset 160..224 winning hash, 64-byte BE (out, only when found)
//! ```
//!
//! JS calls `solve(iters)` in chunks (e.g. 5–20k iterations) to keep the
//! worker responsive; a return of 0 means "keep going" (the nonce field has
//! been advanced), 1 means the nonce field now holds a valid share.
use primitive_types::U512;
const HEADER_OFF: usize = 0;
const NONCE_OFF: usize = 32;
const TARGET_OFF: usize = 96;
const HASH_OFF: usize = 160;
const IO_SIZE: usize = 224;
static mut IO: [u8; IO_SIZE] = [0u8; IO_SIZE];
/// Pointer to the shared I/O buffer in WASM linear memory.
#[no_mangle]
pub extern "C" fn io_ptr() -> *mut u8 {
core::ptr::addr_of_mut!(IO) as *mut u8
}
/// Size of the shared I/O buffer, for JS-side sanity checks.
#[no_mangle]
pub extern "C" fn io_len() -> u32 {
IO_SIZE as u32
}
/// Grind up to `iters` nonces. Returns 1 if a share was found.
#[no_mangle]
pub extern "C" fn solve(iters: u32) -> u32 {
let io = unsafe { &mut *core::ptr::addr_of_mut!(IO) };
let mut header = [0u8; 32];
header.copy_from_slice(&io[HEADER_OFF..HEADER_OFF + 32]);
let mut nonce = U512::from_big_endian(&io[NONCE_OFF..NONCE_OFF + 64]);
let target = U512::from_big_endian(&io[TARGET_OFF..TARGET_OFF + 64]);
for _ in 0..iters {
let nonce_bytes = nonce.to_big_endian();
let hash = qpow_math::get_nonce_hash(header, nonce_bytes);
if hash < target {
io[NONCE_OFF..NONCE_OFF + 64].copy_from_slice(&nonce_bytes);
io[HASH_OFF..HASH_OFF + 64].copy_from_slice(&hash.to_big_endian());
return 1;
}
nonce = nonce.saturating_add(U512::one());
}
// Not found: persist progress so the next chunk resumes where we stopped.
io[NONCE_OFF..NONCE_OFF + 64].copy_from_slice(&nonce.to_big_endian());
0
}
#[cfg(test)]
mod tests {
use super::*;
#[test]
fn solve_matches_pool_verification() {
// Difficulty 4 -> target = MAX/4; expect a solution within a few iters.
let header = [7u8; 32];
let difficulty = U512::from(4u64);
let target = U512::MAX / difficulty;
let io = unsafe { &mut *core::ptr::addr_of_mut!(IO) };
io[HEADER_OFF..HEADER_OFF + 32].copy_from_slice(&header);
let start = U512::from(1u64) << 64; // arbitrary "session range" start
io[NONCE_OFF..NONCE_OFF + 64].copy_from_slice(&start.to_big_endian());
io[TARGET_OFF..TARGET_OFF + 64].copy_from_slice(&target.to_big_endian());
let mut found = 0;
for _ in 0..64 {
found = solve(256);
if found == 1 {
break;
}
}
assert_eq!(found, 1, "should find a share at difficulty 4");
let nonce = U512::from_big_endian(&io[NONCE_OFF..NONCE_OFF + 64]);
let (valid, hash) = qpow_math::is_valid_nonce(header, nonce.to_big_endian(), difficulty);
assert!(
valid,
"pool-side verification must accept the solver's nonce"
);
let reported = U512::from_big_endian(&io[HASH_OFF..HASH_OFF + 64]);
assert_eq!(hash, reported);
}
}
examples/systemd/overrides/10-shared-hardware.conf
+4 -5
View File
@@ -14,8 +14,8 @@
#
# Notes:
# - CPUAffinity must be a subset of the service’s allowed cpuset (e.g., cgroups).
# - The miner logs the detected cpuset mask at startup and, when metrics are enabled,
# exports a miner_effective_cpus gauge for dashboards.
# - The miner counts CPUs via the process affinity mask and exports a
# miner_effective_cpus gauge on the always-on metrics endpoint.
# - Adjust the CPU list and worker count to match your hardware and cpuset.
#
# Example below assumes an 8‑way logical CPU system (IDs 0‑7).
@@ -40,9 +40,8 @@ IOSchedulingPriority=4
# Default time-sharing policy is fine for CPU-bound workloads when pinned.
CPUSchedulingPolicy=other
# Optional: pin worker count and enable metrics (uncomment and adjust as needed).
# Environment=MINER_ENGINE=cpu-fast
# Environment=MINER_WORKERS=4
# Optional: pin worker count and change the metrics port (uncomment and adjust as needed).
# Environment=MINER_CPU_WORKERS=4
# Environment=MINER_METRICS_PORT=9900
# Environment=RUST_LOG=info,miner=debug
examples/systemd/overrides/20-dedicated-hardware.conf
+10 -11
View File
@@ -15,11 +15,11 @@
# - CPUAffinity MUST be a subset of the service’s allowed cpuset (cgroup v2:
# /sys/fs/cgroup/cpuset.cpus.effective). If your deployment already constrains
# the miner with a cpuset, you can omit CPUAffinity here to inherit that mask.
# - The miner logs the detected cpuset mask at startup (debug) and, when metrics
# are enabled, exports miner_effective_cpus for dashboards.
# - Set MINER_WORKERS to match the number of CPUs in your affinity mask if you want
# full utilization; otherwise the miner default (50% of effective CPUs) is used
# when --workers is omitted.
# - The miner counts CPUs via the process affinity mask and exports
# miner_effective_cpus on the always-on metrics endpoint.
# - Set MINER_CPU_WORKERS to match the number of CPUs in your affinity mask if you
# want full utilization; otherwise the miner auto-detects when --cpu-workers is
# omitted.
[Service]
# Bind miner to ALL logical CPUs intended for mining. EDIT this to match your host.
@@ -46,10 +46,9 @@ IOSchedulingPriority=0
TasksMax=infinity
MemoryMax=infinity
# Optional: pin worker count and enable metrics (uncomment and adjust).
# Ensure MINER_WORKERS matches the size of CPUAffinity (or the cpuset mask) to use all CPUs.
# Environment=MINER_ENGINE=cpu-fast
# Environment=MINER_WORKERS=16
# Optional: pin worker count and change the metrics port (uncomment and adjust).
# Ensure MINER_CPU_WORKERS matches the size of CPUAffinity (or the cpuset mask) to use all CPUs.
# Environment=MINER_CPU_WORKERS=16
# Environment=MINER_METRICS_PORT=9900
# Environment=RUST_LOG=info,miner=debug
@@ -64,5 +63,5 @@ MemoryMax=infinity
# and compare with:
# cat /sys/fs/cgroup/cpuset.cpus.effective
#
# - The service logs the cpuset mask at startup and the metrics exporter (if enabled)
# exposes miner_effective_cpus for observability.
# - The always-on metrics exporter exposes miner_effective_cpus for
# observability of the detected CPU count (MINER_METRICS_PORT changes the port).
examples/systemd/quantus-miner.service
+19 -7
View File
@@ -15,14 +15,24 @@
#
# Notes:
# - The miner CLI supports environment variables for all flags (see README):
# MINER_PORT, MINER_METRICS_PORT, MINER_WORKERS, MINER_ENGINE,
# MINER_PROGRESS_CHUNK_MS,
# MINER_MANIP_SOLVED_BLOCKS, MINER_MANIP_BASE_DELAY_NS,
# MINER_MANIP_STEP_BATCH, MINER_MANIP_THROTTLE_CAP.
# MINER_NODE_ADDR, MINER_AUTH_TOKEN_FILE (or MINER_AUTH_TOKEN),
# MINER_TLS_CERT_SHA256_FILE (or MINER_TLS_CERT_SHA256),
# MINER_CPU_WORKERS, MINER_GPU_DEVICES, MINER_GPU_BATCH_SIZE,
# MINER_CPU_BATCH_SIZE, MINER_GPU_THROTTLE_MS, MINER_METRICS_PORT,
# MINER_ALLOW_INTEGRATED, MINER_VERBOSE.
# This unit relies on those env vars (set in /etc/default|/etc/sysconfig) so
# ExecStart can remain stable and simple.
# - The service logs the detected cpuset mask (if any) and emits a
# miner_effective_cpus metric when metrics are enabled.
# - The node requires miner auth: MINER_AUTH_TOKEN_FILE and
# MINER_TLS_CERT_SHA256_FILE (or their inline variants) are mandatory —
# the miner exits at startup without them. Copy the node's
# miner-auth-token and miner-tls-cert-sha256 files somewhere this service
# can read (e.g. /etc/quantus-miner/): ProtectHome=true below blocks
# paths under /home and /root, where the node's base path usually lives.
# - The miner counts CPUs via the process affinity mask (so CPUAffinity=
# is respected) and emits a miner_effective_cpus metric.
# - The Prometheus exporter is ALWAYS on and binds plaintext HTTP on
# 0.0.0.0:9900 by default; MINER_METRICS_PORT only changes the port.
# Firewall the port if the host is reachable from untrusted networks.
[Unit]
Description=Quantus External Miner Service
@@ -53,7 +63,9 @@ Environment=RUST_LOG=info
# Launch the miner. All configuration should be passed via environment vars
# (preferred) or via EXTRA_MINER_FLAGS in the env file for additional arguments.
ExecStart=/usr/local/bin/quantus-miner $EXTRA_MINER_FLAGS
# MINER_AUTH_TOKEN_FILE and MINER_TLS_CERT_SHA256_FILE must be set in the env
# file (see readme.md) or the miner exits immediately.
ExecStart=/usr/local/bin/quantus-miner serve $EXTRA_MINER_FLAGS
# Graceful shutdown and restart behavior
Restart=always
examples/systemd/readme.md
+73 -36
View File
@@ -18,6 +18,18 @@ Prerequisites
- The quantus-miner binary installed at /usr/local/bin/quantus-miner (or adjust ExecStart).
- A service account (recommended):
- sudo useradd --system --no-create-home --shell /usr/sbin/nologin quantus
- The node's miner auth token and TLS cert fingerprint (required since the node
authenticates miners). The node creates both on first start with
--miner-listen-port and logs their paths (default:
`<base-path>/chains/<chain>/miner-auth-token` and
`<base-path>/chains/<chain>/miner-tls-cert-sha256`). Copy them where the
service can read them — the unit sets ProtectHome=true, so paths under /home
or /root are not readable:
- sudo install -d -m 0755 /etc/quantus-miner
- sudo install -m 0640 -o root -g quantus \
"<base-path>/chains/<chain>/miner-auth-token" /etc/quantus-miner/miner-auth-token
- sudo install -m 0644 \
"<base-path>/chains/<chain>/miner-tls-cert-sha256" /etc/quantus-miner/miner-tls-cert-sha256
Install (unit)
1) Copy the service file
@@ -26,21 +38,22 @@ Install (unit)
2) Create a writable working directory (managed by systemd via StateDirectory)
sudo install -d -o quantus -g quantus /var/lib/quantus-miner
3) (Optional) Provide environment variables
3) Provide environment variables (auth vars are REQUIRED)
- Debian/Ubuntu: sudoedit /etc/default/quantus-miner
- RHEL/CentOS/Fed: sudoedit /etc/sysconfig/quantus-miner
Common variables (examples):
MINER_ENGINE=cpu
MINER_PORT=9833
MINER_METRICS_PORT=9900 # enable Prometheus exporter
MINER_WORKERS=4 # leave unset to use default (50% of effective CPUs)
MINER_PROGRESS_CHUNK_MS=2000
# Throttling engine (cpu-chain-manipulator) knobs:
# MINER_MANIP_SOLVED_BLOCKS=0
# MINER_MANIP_BASE_DELAY_NS=500000
# MINER_MANIP_STEP_BATCH=10000
# MINER_MANIP_THROTTLE_CAP=0
Common variables (examples). Note: systemd EnvironmentFile= does not strip
inline comments — keep each assignment on its own line with nothing after
the value.
MINER_NODE_ADDR=127.0.0.1:9833
# Required: node auth token + TLS cert pin (see Prerequisites)
MINER_AUTH_TOKEN_FILE=/etc/quantus-miner/miner-auth-token
MINER_TLS_CERT_SHA256_FILE=/etc/quantus-miner/miner-tls-cert-sha256
# CPU worker threads; leave unset to auto-detect (~50% of available CPUs)
MINER_CPU_WORKERS=4
MINER_GPU_DEVICES=0
# Prometheus exporter port (the exporter is always on; this only changes the port)
MINER_METRICS_PORT=9900
# Extra CLI flags (kept stable ExecStart):
# EXTRA_MINER_FLAGS="--some-future-flag value"
@@ -65,31 +78,44 @@ Dedicated hardware (miner only)
sudo systemctl restart quantus-miner
Configuration reference (environment variables)
- MINER_ENGINE
- cpu (default), cpu-chain-manipulator
- gpu for high-performance GPU mining
- MINER_PORT
- HTTP API port (default 9833)
- MINER_NODE_ADDR
- Address of the node's miner QUIC endpoint (default 127.0.0.1:9833).
- MINER_AUTH_TOKEN_FILE (required, or MINER_AUTH_TOKEN inline)
- Path to a copy of the node's miner-auth-token file. The file variant is
preferred so the secret stays off the command line and out of `systemctl show`.
- MINER_TLS_CERT_SHA256_FILE (required, or MINER_TLS_CERT_SHA256 inline)
- Path to a copy of the node's miner-tls-cert-sha256 file (the fingerprint is
also printed in the node's startup logs).
- MINER_CPU_WORKERS
- CPU worker threads. If unset, auto-detected.
- MINER_GPU_DEVICES
- Number of GPU devices to use. If unset, auto-detected.
- MINER_GPU_BATCH_SIZE / MINER_CPU_BATCH_SIZE
- Nonces/hashes per cancellation check (defaults 1000000 / 10000).
- MINER_GPU_THROTTLE_MS
- Delay between GPU batches in milliseconds (default 0 = no throttle).
- MINER_METRICS_PORT
- Enable Prometheus exporter when set (e.g., 9900). If unset, metrics exporter is disabled.
- MINER_WORKERS
- Worker threads (logical CPUs). If unset, defaults to ~50% of effective CPUs (clamped to [1, effective-1]).
- MINER_PROGRESS_CHUNK_MS
- Target milliseconds for per-thread progress updates (default 2000ms).
- Throttling engine (cpu-chain-manipulator) knobs
- MINER_MANIP_SOLVED_BLOCKS, MINER_MANIP_BASE_DELAY_NS, MINER_MANIP_STEP_BATCH, MINER_MANIP_THROTTLE_CAP
- Prometheus exporter port (default 9900). The exporter is ALWAYS on and
binds plaintext HTTP on all interfaces (0.0.0.0); this variable only
changes the port — there is no disable or loopback-only option. Firewall
the port or restrict it to your monitoring network.
- MINER_ALLOW_INTEGRATED
- Allow integrated GPUs even when discrete GPUs are present.
- EXTRA_MINER_FLAGS
- Optional extra CLI flags appended to ExecStart.
CPU affinity, cpusets, and workers
- The miner detects the effective CPU capacity (logical CPUs) visible to the process by preferring cgroup v2 cpuset (cpuset.cpus.effective), falling back to v1, else using all logical CPUs.
- At startup (debug level), the miner logs the detected cpuset mask (if any).
- A Prometheus gauge miner_effective_cpus is emitted (when metrics are enabled) with the effective count for dashboards/alerts.
- If --workers (or MINER_WORKERS) exceeds effective CPUs, it is clamped and a warning is logged.
- If omitted, the miner defaults to ~50% of effective CPUs (but always at least 1 and less than or equal to effective-1).
- The miner counts available CPUs via the process CPU affinity mask (num_cpus),
so a systemd CPUAffinity= setting (or a cgroup cpuset) is reflected in the count.
- If MINER_CPU_WORKERS is unset, the miner uses ~50% of the detected CPUs
(at least 1) and logs the choice at startup ("Auto-detected N CPU workers").
- An explicit MINER_CPU_WORKERS value is used as-is — it is NOT clamped to the
affinity mask, so an oversized value oversubscribes the pinned CPUs.
- A Prometheus gauge miner_effective_cpus is emitted with the detected count for dashboards/alerts.
- When pinning CPUAffinity at the systemd level:
- Ensure CPUAffinity is a subset of the cgroup cpuset mask.
- Consider setting MINER_WORKERS to match the number of CPUs in the affinity mask if you want full utilization, or rely on the default 50% policy.
- Set MINER_CPU_WORKERS to the number of CPUs in the affinity mask for full
utilization, or leave it unset for the ~50% default.
Security hardening (in the unit)
- NoNewPrivileges=true
@@ -101,6 +127,9 @@ Security hardening (in the unit)
- RestrictSUIDSGID=true
- SystemCallFilter=@system-service
Adjust or relax as needed for your environment.
Note: the Prometheus exporter always listens on 0.0.0.0:<MINER_METRICS_PORT>
(default 9900) — the unit cannot disable it or bind it to loopback. Firewall
the port if the host is reachable from untrusted networks.
Validation and troubleshooting
- Check service status and logs:
@@ -110,20 +139,28 @@ Validation and troubleshooting
taskset -cp "$pid"
- Verify cpuset mask (cgroup v2):
cat /sys/fs/cgroup/cpuset.cpus.effective
- Metrics:
- If MINER_METRICS_PORT is set, curl http://127.0.0.1:<port>/metrics
- Metrics (always on, default port 9900):
- curl http://127.0.0.1:${MINER_METRICS_PORT:-9900}/metrics
- Look for miner_effective_cpus and per-job/thread metrics.
- Common pitfalls:
- ExecStart path wrong (ensure /usr/local/bin/quantus-miner exists and is executable).
- Service user/group missing (create quantus or adjust User/Group).
- Miner exits immediately with "miner auth token required" / "TLS cert fingerprint required":
set MINER_AUTH_TOKEN_FILE and MINER_TLS_CERT_SHA256_FILE in the env file (see Prerequisites).
- Auth/TLS files unreadable: ProtectHome=true blocks /home and /root; copy the
files to /etc/quantus-miner/ and make them readable by the quantus user.
- Node rejects the miner after the node's base path changed or its credential
files (miner-auth-token, miner-tls-cert.der/-key.der) were deleted and
regenerated: re-copy both files and restart. (A plain purge-chain does NOT
rotate them — it removes only the database, so the same token and cert are
reloaded on the next start.)
- CPUAffinity not a subset of the cgroup cpuset (adjust cpuset or affinity).
- MINER_ENGINE set to gpu-* (currently unimplemented, exits with clear error).
- Insufficient permissions to write WorkingDirectory (systemd StateDirectory creates /var/lib/quantus-miner with correct ownership).
Operational tips
- For shared machines: prefer 10-shared-hardware.conf and leave MINER_WORKERS unset (defaults to ~50%).
- For dedicated machines: use 20-dedicated-hardware.conf and set MINER_WORKERS to the number of CPUs in CPUAffinity (or omit CPUAffinity to inherit cpuset).
- Use RUST_LOG=info,miner=debug temporarily to verify startup detection (cpuset mask, effective CPUs) and to observe mining loop behavior; then turn back down to reduce log volume.
- For shared machines: prefer 10-shared-hardware.conf and leave MINER_CPU_WORKERS unset (auto-detect).
- For dedicated machines: use 20-dedicated-hardware.conf and set MINER_CPU_WORKERS to the number of CPUs in CPUAffinity (or omit CPUAffinity to inherit cpuset).
- Use RUST_LOG=info,miner=debug temporarily to verify startup detection (worker auto-detection, GPU discovery) and to observe mining loop behavior; then turn back down to reduce log volume.
Support
- Repository: https://github.com/Quantus-Network/quantus-miner
rust-toolchain
+1 -1
View File
@@ -1,4 +1,4 @@
[toolchain]
channel = "stable"
channel = "1.93.0"
components = ["clippy", "rustfmt"]
profile = "minimal"