Compare commits
17 Commits
| Author | SHA1 | Date | |
|---|---|---|---|
|
aa9e7ca5f2 | ||
|
|
6902c650a9 | ||
|
|
e0f7beef74 | ||
|
|
e92c4f0769 | ||
|
|
06ddee497d | ||
|
|
63d7e114c2 | ||
|
|
57f69b6aae | ||
|
|
f35660c9e9 | ||
|
|
3466eff8b0 | ||
|
|
c0bfca592d | ||
|
|
6912470e84 | ||
|
|
534b3ac2f8 | ||
|
|
ff3e65361a | ||
|
|
a6660c5e13 | ||
|
|
85006dd3a7 | ||
|
|
0a96d2f6b7 | ||
|
|
aa1204bab9 |
146
.github/workflows/cuda-builder.yml
vendored
146
.github/workflows/cuda-builder.yml
vendored
@@ -1,83 +1,83 @@
|
||||
name: Cuda Docker Image Builder
|
||||
# name: Cuda Docker Image Builder
|
||||
|
||||
on:
|
||||
workflow_dispatch:
|
||||
push:
|
||||
branches:
|
||||
- main
|
||||
paths:
|
||||
- cuda-builder.cf
|
||||
pull_request:
|
||||
branches:
|
||||
- main
|
||||
paths:
|
||||
- cuda-builder.cf
|
||||
# on:
|
||||
# workflow_dispatch:
|
||||
# push:
|
||||
# branches:
|
||||
# - main
|
||||
# paths:
|
||||
# - cuda-builder.cf
|
||||
# pull_request:
|
||||
# branches:
|
||||
# - main
|
||||
# paths:
|
||||
# - cuda-builder.cf
|
||||
|
||||
permissions:
|
||||
contents: read
|
||||
packages: write
|
||||
# permissions:
|
||||
# contents: read
|
||||
# packages: write
|
||||
|
||||
jobs:
|
||||
build-and-push:
|
||||
name: cuda-builder:${{ matrix.cuda_tag }}
|
||||
runs-on: ubuntu-latest
|
||||
strategy:
|
||||
fail-fast: false
|
||||
matrix:
|
||||
cuda_tag:
|
||||
- 12.9.0
|
||||
- 13.0.0
|
||||
steps:
|
||||
- shell: bash
|
||||
run: |
|
||||
echo "IMAGE_NAME=ghcr.io/${GITHUB_REPOSITORY_OWNER,,}/cuda-builder" >> "$GITHUB_ENV"
|
||||
# jobs:
|
||||
# build-and-push:
|
||||
# name: cuda-builder:${{ matrix.cuda_tag }}
|
||||
# runs-on: ubuntu-latest
|
||||
# strategy:
|
||||
# fail-fast: false
|
||||
# matrix:
|
||||
# cuda_tag:
|
||||
# - 12.9.0
|
||||
# - 13.0.0
|
||||
# steps:
|
||||
# - shell: bash
|
||||
# run: |
|
||||
# echo "IMAGE_NAME=ghcr.io/${GITHUB_REPOSITORY_OWNER,,}/cuda-builder" >> "$GITHUB_ENV"
|
||||
|
||||
- name: Checkout repository
|
||||
uses: actions/checkout@v4
|
||||
# - name: Checkout repository
|
||||
# uses: actions/checkout@v4
|
||||
|
||||
- uses: ./.github/actions/disk
|
||||
# - uses: ./.github/actions/disk
|
||||
|
||||
- name: Set up Docker Buildx
|
||||
uses: docker/setup-buildx-action@v3
|
||||
# - name: Set up Docker Buildx
|
||||
# uses: docker/setup-buildx-action@v3
|
||||
|
||||
- name: Prepare Buildx cache
|
||||
uses: actions/cache@v4
|
||||
with:
|
||||
path: /tmp/.buildx-cache
|
||||
key: ${{ runner.os }}-buildx-cuda-builder-${{ matrix.cuda_tag }}-${{ github.sha }}
|
||||
restore-keys: |
|
||||
${{ runner.os }}-buildx-cuda-builder-${{ matrix.cuda_tag }}-
|
||||
# - name: Prepare Buildx cache
|
||||
# uses: actions/cache@v4
|
||||
# with:
|
||||
# path: /tmp/.buildx-cache
|
||||
# key: ${{ runner.os }}-buildx-cuda-builder-${{ matrix.cuda_tag }}-${{ github.sha }}
|
||||
# restore-keys: |
|
||||
# ${{ runner.os }}-buildx-cuda-builder-${{ matrix.cuda_tag }}-
|
||||
|
||||
- name: Log in to GHCR
|
||||
if: github.event_name == 'push'
|
||||
uses: docker/login-action@v3
|
||||
with:
|
||||
registry: ghcr.io
|
||||
username: ${{ github.repository_owner }}
|
||||
password: ${{ secrets.GITHUB_TOKEN }}
|
||||
# - name: Log in to GHCR
|
||||
# if: github.event_name == 'push'
|
||||
# uses: docker/login-action@v3
|
||||
# with:
|
||||
# registry: ghcr.io
|
||||
# username: ${{ github.repository_owner }}
|
||||
# password: ${{ secrets.GITHUB_TOKEN }}
|
||||
|
||||
- name: Build and (conditionally) push cuda-builder image
|
||||
uses: docker/build-push-action@v5
|
||||
with:
|
||||
context: .
|
||||
file: cuda-builder.cf
|
||||
build-args: |
|
||||
CUDA_TAG=${{ matrix.cuda_tag }}
|
||||
pull: true
|
||||
push: ${{ github.event_name == 'push' }}
|
||||
tags: |
|
||||
${{ env.IMAGE_NAME }}:${{ matrix.cuda_tag }}
|
||||
labels: |
|
||||
org.opencontainers.image.source=${{ github.repository }}
|
||||
org.opencontainers.image.description=Pristine CUDA builder image (CUDA ${{ matrix.cuda_tag }})
|
||||
org.opencontainers.image.licenses=Apache-2.0
|
||||
org.opencontainers.image.title=cuda-builder
|
||||
org.opencontainers.image.vendor=Quantus Network
|
||||
cache-from: type=local,src=/tmp/.buildx-cache
|
||||
cache-to: type=local,dest=/tmp/.buildx-cache-new,mode=max
|
||||
provenance: false
|
||||
# - name: Build and (conditionally) push cuda-builder image
|
||||
# uses: docker/build-push-action@v5
|
||||
# with:
|
||||
# context: .
|
||||
# file: cuda-builder.cf
|
||||
# build-args: |
|
||||
# CUDA_TAG=${{ matrix.cuda_tag }}
|
||||
# pull: true
|
||||
# push: ${{ github.event_name == 'push' }}
|
||||
# tags: |
|
||||
# ${{ env.IMAGE_NAME }}:${{ matrix.cuda_tag }}
|
||||
# labels: |
|
||||
# org.opencontainers.image.source=${{ github.repository }}
|
||||
# org.opencontainers.image.description=Pristine CUDA builder image (CUDA ${{ matrix.cuda_tag }})
|
||||
# org.opencontainers.image.licenses=Apache-2.0
|
||||
# org.opencontainers.image.title=cuda-builder
|
||||
# org.opencontainers.image.vendor=Quantus Network
|
||||
# cache-from: type=local,src=/tmp/.buildx-cache
|
||||
# cache-to: type=local,dest=/tmp/.buildx-cache-new,mode=max
|
||||
# provenance: false
|
||||
|
||||
- name: Move Buildx cache
|
||||
run: |
|
||||
rm -rf /tmp/.buildx-cache
|
||||
mv /tmp/.buildx-cache-new /tmp/.buildx-cache
|
||||
# - name: Move Buildx cache
|
||||
# run: |
|
||||
# rm -rf /tmp/.buildx-cache
|
||||
# mv /tmp/.buildx-cache-new /tmp/.buildx-cache
|
||||
|
||||
13
.github/workflows/release-proposal.yml
vendored
13
.github/workflows/release-proposal.yml
vendored
@@ -96,15 +96,9 @@ jobs:
|
||||
|
||||
# Increment based on type
|
||||
if [[ "$VERSION_TYPE" == "major" ]]; then
|
||||
if [[ "$major" == "0" ]]; then # Handle 0.x.y -> 0.(x+1).0
|
||||
major=$major
|
||||
minor=$((minor + 1))
|
||||
patch=0
|
||||
else
|
||||
major=$((major + 1))
|
||||
minor=0
|
||||
patch=0
|
||||
fi
|
||||
major=$((major + 1))
|
||||
minor=0
|
||||
patch=0
|
||||
elif [[ "$VERSION_TYPE" == "minor" ]]; then
|
||||
minor=$((minor + 1))
|
||||
patch=0
|
||||
@@ -166,7 +160,6 @@ jobs:
|
||||
cargo update -p engine-cpu --precise "$new_cargo_version"
|
||||
cargo update -p engine-gpu-cuda --precise "$new_cargo_version"
|
||||
cargo update -p engine-gpu-opencl --precise "$new_cargo_version"
|
||||
cargo update -p engine-montgomery --precise "$new_cargo_version"
|
||||
|
||||
# Verify everything compiles correctly
|
||||
cargo check --workspace
|
||||
|
||||
472
Cargo.lock
generated
472
Cargo.lock
generated
@@ -26,6 +26,12 @@ dependencies = [
|
||||
"memchr",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "anes"
|
||||
version = "0.1.6"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "4b46cbb362ab8752921c97e041f5e366ee6297bd428a31275b9fcf1e380f7299"
|
||||
|
||||
[[package]]
|
||||
name = "anstream"
|
||||
version = "0.6.20"
|
||||
@@ -187,6 +193,12 @@ version = "1.10.1"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "d71b6127be86fdcfddb610f7182ac57211d4b18a3e9c82eb2d17662f2227ad6a"
|
||||
|
||||
[[package]]
|
||||
name = "cast"
|
||||
version = "0.3.0"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "37b2a672a2cb129a2e41c10b1224bb368f9f37a2b16b612598138befd7b37eb5"
|
||||
|
||||
[[package]]
|
||||
name = "cc"
|
||||
version = "1.2.38"
|
||||
@@ -203,6 +215,33 @@ version = "1.0.3"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "2fd1289c04a9ea8cb22300a459a72a385d7c73d3259e2ed7dcb2af674838cfa9"
|
||||
|
||||
[[package]]
|
||||
name = "ciborium"
|
||||
version = "0.2.2"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "42e69ffd6f0917f5c029256a24d0161db17cea3997d185db0d35926308770f0e"
|
||||
dependencies = [
|
||||
"ciborium-io",
|
||||
"ciborium-ll",
|
||||
"serde",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "ciborium-io"
|
||||
version = "0.2.2"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "05afea1e0a06c9be33d539b876f1ce3692f4afea2cb41f740e7743225ed1c757"
|
||||
|
||||
[[package]]
|
||||
name = "ciborium-ll"
|
||||
version = "0.2.2"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "57663b653d948a338bfb3eeba9bb2fd5fcfaecb9e199e87e1eda4d9e8b240fd9"
|
||||
dependencies = [
|
||||
"ciborium-io",
|
||||
"half",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "cl-sys"
|
||||
version = "0.4.3"
|
||||
@@ -287,6 +326,42 @@ dependencies = [
|
||||
"libc",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "criterion"
|
||||
version = "0.5.1"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "f2b12d017a929603d80db1831cd3a24082f8137ce19c69e6447f54f5fc8d692f"
|
||||
dependencies = [
|
||||
"anes",
|
||||
"cast",
|
||||
"ciborium",
|
||||
"clap",
|
||||
"criterion-plot",
|
||||
"is-terminal",
|
||||
"itertools 0.10.5",
|
||||
"num-traits 0.2.19",
|
||||
"once_cell",
|
||||
"oorandom",
|
||||
"plotters",
|
||||
"rayon",
|
||||
"regex",
|
||||
"serde",
|
||||
"serde_derive",
|
||||
"serde_json",
|
||||
"tinytemplate",
|
||||
"walkdir",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "criterion-plot"
|
||||
version = "0.5.0"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "6b50826342786a51a89e2da3a28f1c32b06e387201bc2d19791f622c673706b1"
|
||||
dependencies = [
|
||||
"cast",
|
||||
"itertools 0.10.5",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "crossbeam"
|
||||
version = "0.8.4"
|
||||
@@ -349,16 +424,6 @@ version = "0.2.4"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "460fbee9c2c2f33933d720630a6a0bac33ba7053db5344fac858d4b8952d77d5"
|
||||
|
||||
[[package]]
|
||||
name = "crypto-bigint"
|
||||
version = "0.5.5"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "0dc92fb57ca44df6db8059111ab3af99a63d5d0f8375d9972e319a379c6bab76"
|
||||
dependencies = [
|
||||
"rand_core 0.6.4",
|
||||
"subtle",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "crypto-common"
|
||||
version = "0.1.6"
|
||||
@@ -460,6 +525,12 @@ dependencies = [
|
||||
"syn 2.0.106",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "either"
|
||||
version = "1.15.0"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "48c757948c5ede0e46177b7add2e67155f70e33c07fea8284df6576da70b3719"
|
||||
|
||||
[[package]]
|
||||
name = "encoding_rs"
|
||||
version = "0.8.35"
|
||||
@@ -471,19 +542,21 @@ dependencies = [
|
||||
|
||||
[[package]]
|
||||
name = "engine-cpu"
|
||||
version = "0.2.1"
|
||||
version = "1.0.0"
|
||||
dependencies = [
|
||||
"anyhow",
|
||||
"criterion",
|
||||
"hex",
|
||||
"log",
|
||||
"pow-core",
|
||||
"primitive-types",
|
||||
"rand 0.8.5",
|
||||
"thiserror",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "engine-gpu-cuda"
|
||||
version = "0.2.1"
|
||||
version = "1.0.0"
|
||||
dependencies = [
|
||||
"anyhow",
|
||||
"cust",
|
||||
@@ -495,14 +568,14 @@ dependencies = [
|
||||
"num-traits 0.2.19",
|
||||
"pow-core",
|
||||
"primitive-types",
|
||||
"qp-poseidon-core",
|
||||
"rustacuda",
|
||||
"sha3",
|
||||
"thiserror",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "engine-gpu-opencl"
|
||||
version = "0.2.1"
|
||||
version = "1.0.0"
|
||||
dependencies = [
|
||||
"anyhow",
|
||||
"log",
|
||||
@@ -512,20 +585,6 @@ dependencies = [
|
||||
"thiserror",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "engine-montgomery"
|
||||
version = "0.2.1"
|
||||
dependencies = [
|
||||
"crypto-bigint",
|
||||
"engine-cpu",
|
||||
"log",
|
||||
"metrics",
|
||||
"num-traits 0.2.19",
|
||||
"pow-core",
|
||||
"primitive-types",
|
||||
"sha3",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "enum_primitive"
|
||||
version = "0.1.1"
|
||||
@@ -770,6 +829,17 @@ dependencies = [
|
||||
"tracing",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "half"
|
||||
version = "2.7.0"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "e54c115d4f30f52c67202f079c5f9d8b49db4691f460fdb0b4c2e838261b2ba5"
|
||||
dependencies = [
|
||||
"cfg-if",
|
||||
"crunchy",
|
||||
"zerocopy",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "hashbrown"
|
||||
version = "0.16.0"
|
||||
@@ -1058,6 +1128,24 @@ version = "1.70.1"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "7943c866cc5cd64cbc25b2e01621d07fa8eb2a1a23160ee81ce38704e97b8ecf"
|
||||
|
||||
[[package]]
|
||||
name = "itertools"
|
||||
version = "0.10.5"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "b0fd2260e829bddf4cb6ea802289de2f86d6a7a690192fbe91b3f46e0f2c8473"
|
||||
dependencies = [
|
||||
"either",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "itertools"
|
||||
version = "0.14.0"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "2b192c782037fadd9cfa75548310488aabdbf3d2da73885b31bd0abd03351285"
|
||||
dependencies = [
|
||||
"either",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "itoa"
|
||||
version = "1.0.15"
|
||||
@@ -1074,15 +1162,6 @@ dependencies = [
|
||||
"wasm-bindgen",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "keccak"
|
||||
version = "0.1.5"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "ecc2af9a1119c51f12a14607e783cb977bde58bc069ff0c3da1095e635d70654"
|
||||
dependencies = [
|
||||
"cpufeatures",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "lazy_static"
|
||||
version = "1.5.0"
|
||||
@@ -1131,7 +1210,7 @@ checksum = "32a282da65faaf38286cf3be983213fcf1d2e2a58700e808f83f4ea9a4804bc0"
|
||||
|
||||
[[package]]
|
||||
name = "metrics"
|
||||
version = "0.2.1"
|
||||
version = "1.0.0"
|
||||
dependencies = [
|
||||
"anyhow",
|
||||
"log",
|
||||
@@ -1161,7 +1240,7 @@ dependencies = [
|
||||
|
||||
[[package]]
|
||||
name = "miner-cli"
|
||||
version = "0.2.1"
|
||||
version = "1.0.0"
|
||||
dependencies = [
|
||||
"clap",
|
||||
"env_logger",
|
||||
@@ -1172,14 +1251,13 @@ dependencies = [
|
||||
|
||||
[[package]]
|
||||
name = "miner-service"
|
||||
version = "0.2.1"
|
||||
version = "1.0.0"
|
||||
dependencies = [
|
||||
"anyhow",
|
||||
"crossbeam-channel",
|
||||
"engine-cpu",
|
||||
"engine-gpu-cuda",
|
||||
"engine-gpu-opencl",
|
||||
"engine-montgomery",
|
||||
"hex",
|
||||
"log",
|
||||
"metrics",
|
||||
@@ -1197,7 +1275,7 @@ dependencies = [
|
||||
|
||||
[[package]]
|
||||
name = "miner-telemetry"
|
||||
version = "0.2.1"
|
||||
version = "1.0.0"
|
||||
dependencies = [
|
||||
"anyhow",
|
||||
"futures 0.3.31",
|
||||
@@ -1378,6 +1456,128 @@ version = "1.70.1"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "a4895175b425cb1f87721b59f0f286c2092bd4af812243672510e1ac53e2e0ad"
|
||||
|
||||
[[package]]
|
||||
name = "oorandom"
|
||||
version = "11.1.5"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "d6790f58c7ff633d8771f42965289203411a5e5c68388703c06e14f24770b41e"
|
||||
|
||||
[[package]]
|
||||
name = "p3-dft"
|
||||
version = "0.3.0"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "e3b2764a3982d22d62aa933c8de6f9d71d8a474c9110b69e675dea1887bdeffc"
|
||||
dependencies = [
|
||||
"itertools 0.14.0",
|
||||
"p3-field",
|
||||
"p3-matrix",
|
||||
"p3-maybe-rayon",
|
||||
"p3-util",
|
||||
"tracing",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "p3-field"
|
||||
version = "0.3.0"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "fc13a73509fe09c67b339951ca8d4cc6e61c9bf08c130dbc90dda52452918cc2"
|
||||
dependencies = [
|
||||
"itertools 0.14.0",
|
||||
"num-bigint",
|
||||
"p3-maybe-rayon",
|
||||
"p3-util",
|
||||
"paste",
|
||||
"rand 0.9.2",
|
||||
"serde",
|
||||
"tracing",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "p3-goldilocks"
|
||||
version = "0.3.0"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "552849f6309ffde34af0d31aa9a2d0a549cb0ec138d9792bfbf4a17800742362"
|
||||
dependencies = [
|
||||
"num-bigint",
|
||||
"p3-dft",
|
||||
"p3-field",
|
||||
"p3-mds",
|
||||
"p3-poseidon2",
|
||||
"p3-symmetric",
|
||||
"p3-util",
|
||||
"paste",
|
||||
"rand 0.9.2",
|
||||
"serde",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "p3-matrix"
|
||||
version = "0.3.0"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "c8e1e9f69c2fe15768b3ceb2915edb88c47398aa22c485d8163deab2a47fe194"
|
||||
dependencies = [
|
||||
"itertools 0.14.0",
|
||||
"p3-field",
|
||||
"p3-maybe-rayon",
|
||||
"p3-util",
|
||||
"rand 0.9.2",
|
||||
"serde",
|
||||
"tracing",
|
||||
"transpose",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "p3-maybe-rayon"
|
||||
version = "0.3.0"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "33f765046b763d046728b3246b690f81dfa7ccd7523b7a1582c74f616fbce6a0"
|
||||
|
||||
[[package]]
|
||||
name = "p3-mds"
|
||||
version = "0.3.0"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "6c90541c6056712daf2ee69ec328db8b5605ae8dbafe60226c8eb75eaac0e1f9"
|
||||
dependencies = [
|
||||
"p3-dft",
|
||||
"p3-field",
|
||||
"p3-symmetric",
|
||||
"p3-util",
|
||||
"rand 0.9.2",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "p3-poseidon2"
|
||||
version = "0.3.0"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "88e9f053f120a78ad27e9c1991a0ea547777328ca24025c42364d6ee2667d59a"
|
||||
dependencies = [
|
||||
"p3-field",
|
||||
"p3-mds",
|
||||
"p3-symmetric",
|
||||
"p3-util",
|
||||
"rand 0.9.2",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "p3-symmetric"
|
||||
version = "0.3.0"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "72d5db8f05a26d706dfd8aaf7aa4272ca4f3e7a075db897ec7108f24fad78759"
|
||||
dependencies = [
|
||||
"itertools 0.14.0",
|
||||
"p3-field",
|
||||
"serde",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "p3-util"
|
||||
version = "0.3.0"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "6dfee67245d9ce78a15176728da2280032f0a84b5819a39a953e7ec03cfd9bd7"
|
||||
dependencies = [
|
||||
"serde",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "parity-scale-codec"
|
||||
version = "3.7.5"
|
||||
@@ -1429,6 +1629,12 @@ dependencies = [
|
||||
"windows-targets 0.52.6",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "paste"
|
||||
version = "1.0.15"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "57c0d7b74b563b49d38dae00a0c37d4d6de9b432382b2892f0574ddcae73fd0a"
|
||||
|
||||
[[package]]
|
||||
name = "percent-encoding"
|
||||
version = "2.3.2"
|
||||
@@ -1467,6 +1673,34 @@ version = "0.1.0"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "8b870d8c151b6f2fb93e84a13146138f05d02ed11c7e7c54f8826aaaf7c9f184"
|
||||
|
||||
[[package]]
|
||||
name = "plotters"
|
||||
version = "0.3.7"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "5aeb6f403d7a4911efb1e33402027fc44f29b5bf6def3effcc22d7bb75f2b747"
|
||||
dependencies = [
|
||||
"num-traits 0.2.19",
|
||||
"plotters-backend",
|
||||
"plotters-svg",
|
||||
"wasm-bindgen",
|
||||
"web-sys",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "plotters-backend"
|
||||
version = "0.3.7"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "df42e13c12958a16b3f7f4386b9ab1f3e7933914ecea48da7139435263a4172a"
|
||||
|
||||
[[package]]
|
||||
name = "plotters-svg"
|
||||
version = "0.3.7"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "51bae2ac328883f7acdfea3d66a7c35751187f870bc81f94563733a154d7a670"
|
||||
dependencies = [
|
||||
"plotters-backend",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "potential_utf"
|
||||
version = "0.1.3"
|
||||
@@ -1478,7 +1712,7 @@ dependencies = [
|
||||
|
||||
[[package]]
|
||||
name = "pow-core"
|
||||
version = "0.2.1"
|
||||
version = "1.0.0"
|
||||
dependencies = [
|
||||
"anyhow",
|
||||
"hex",
|
||||
@@ -1486,8 +1720,8 @@ dependencies = [
|
||||
"num-bigint",
|
||||
"num-traits 0.2.19",
|
||||
"primitive-types",
|
||||
"sha2",
|
||||
"sha3",
|
||||
"qp-poseidon-core",
|
||||
"qpow-math",
|
||||
"thiserror",
|
||||
]
|
||||
|
||||
@@ -1543,6 +1777,46 @@ dependencies = [
|
||||
"thiserror",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "qp-poseidon-constants"
|
||||
version = "1.0.1"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "4d56b56652e9f44a43de9593e75d7c3e0c3a352e10675cf3024e5b3175711cd3"
|
||||
dependencies = [
|
||||
"p3-field",
|
||||
"p3-goldilocks",
|
||||
"p3-poseidon2",
|
||||
"rand 0.9.2",
|
||||
"rand_chacha 0.9.0",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "qp-poseidon-core"
|
||||
version = "1.0.1"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "e658a373a7fb22babeda9ffcc8af0a894e6e3c008272ed735509eccb7769ead3"
|
||||
dependencies = [
|
||||
"p3-field",
|
||||
"p3-goldilocks",
|
||||
"p3-poseidon2",
|
||||
"p3-symmetric",
|
||||
"qp-poseidon-constants",
|
||||
"rand_chacha 0.9.0",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "qpow-math"
|
||||
version = "0.1.0"
|
||||
source = "git+https://github.com/Quantus-Network/chain.git#a5b9e4cc6900641884171a9972e0276ebb73a3ab"
|
||||
dependencies = [
|
||||
"hex",
|
||||
"log",
|
||||
"num-bigint",
|
||||
"num-traits 0.2.19",
|
||||
"primitive-types",
|
||||
"qp-poseidon-core",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "quantus-miner-api"
|
||||
version = "0.0.3"
|
||||
@@ -1642,6 +1916,26 @@ dependencies = [
|
||||
"getrandom 0.3.3",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "rayon"
|
||||
version = "1.11.0"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "368f01d005bf8fd9b1206fb6fa653e6c4a81ceb1466406b81792d87c5677a58f"
|
||||
dependencies = [
|
||||
"either",
|
||||
"rayon-core",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "rayon-core"
|
||||
version = "1.13.0"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "22e18b0f0062d30d4230b2e85ff77fdfe4326feb054b9783a3460d8435c8ab91"
|
||||
dependencies = [
|
||||
"crossbeam-deque",
|
||||
"crossbeam-utils",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "redox_syscall"
|
||||
version = "0.5.17"
|
||||
@@ -1789,6 +2083,15 @@ version = "1.0.20"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "28d3b2b1366ec20994f1fd18c3c594f05c5dd4bc44d8bb0c1c632c8d6829481f"
|
||||
|
||||
[[package]]
|
||||
name = "same-file"
|
||||
version = "1.0.6"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "93fc1dc3aaa9bfed95e02e6eadabb4baf7e3078b0bd1b4d7b6b0b68378900502"
|
||||
dependencies = [
|
||||
"winapi-util",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "scoped-tls"
|
||||
version = "1.0.1"
|
||||
@@ -1873,27 +2176,6 @@ 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",
|
||||
"digest",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "sha3"
|
||||
version = "0.10.8"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "75872d278a8f37ef87fa0ddbda7802605cb18344497949862c0d4dcb291eba60"
|
||||
dependencies = [
|
||||
"digest",
|
||||
"keccak",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "shlex"
|
||||
version = "1.3.0"
|
||||
@@ -1959,6 +2241,12 @@ version = "1.1.0"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "a2eb9349b6444b326872e140eb1cf5e7c522154d69e7a0ffb0fb81c06b37543f"
|
||||
|
||||
[[package]]
|
||||
name = "strength_reduce"
|
||||
version = "0.2.4"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "fe895eb47f22e2ddd4dabc02bce419d2e643c8e3b585c78158b349195bc24d82"
|
||||
|
||||
[[package]]
|
||||
name = "strsim"
|
||||
version = "0.11.1"
|
||||
@@ -2049,6 +2337,16 @@ dependencies = [
|
||||
"zerovec",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "tinytemplate"
|
||||
version = "1.2.1"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "be4d6b5f19ff7664e8c98d03e2139cb510db9b0a60b55f8e8709b689d939b6bc"
|
||||
dependencies = [
|
||||
"serde",
|
||||
"serde_json",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "tokio"
|
||||
version = "1.47.1"
|
||||
@@ -2175,9 +2473,21 @@ checksum = "784e0ac535deb450455cbfa28a6f0df145ea1bb7ae51b821cf5e7927fdcfbdd0"
|
||||
dependencies = [
|
||||
"log",
|
||||
"pin-project-lite",
|
||||
"tracing-attributes",
|
||||
"tracing-core",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "tracing-attributes"
|
||||
version = "0.1.30"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "81383ab64e72a7a8b8e13130c49e3dab29def6d0c7d76a03087b3cf71c5c6903"
|
||||
dependencies = [
|
||||
"proc-macro2",
|
||||
"quote",
|
||||
"syn 2.0.106",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "tracing-core"
|
||||
version = "0.1.34"
|
||||
@@ -2187,6 +2497,16 @@ dependencies = [
|
||||
"once_cell",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "transpose"
|
||||
version = "0.2.3"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "1ad61aed86bc3faea4300c7aee358b4c6d0c8d6ccc36524c96e4c92ccf26e77e"
|
||||
dependencies = [
|
||||
"num-integer",
|
||||
"strength_reduce",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "try-lock"
|
||||
version = "0.2.5"
|
||||
@@ -2334,6 +2654,16 @@ version = "0.9.5"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "0b928f33d975fc6ad9f86c8f283853ad26bdd5b10b7f1542aa2fa15e2289105a"
|
||||
|
||||
[[package]]
|
||||
name = "walkdir"
|
||||
version = "2.5.0"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "29790946404f91d9c5d06f9874efddea1dc06c5efe94541a7d6863108e3a5e4b"
|
||||
dependencies = [
|
||||
"same-file",
|
||||
"winapi-util",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "want"
|
||||
version = "0.3.1"
|
||||
@@ -2455,6 +2785,16 @@ dependencies = [
|
||||
"unicode-ident",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "web-sys"
|
||||
version = "0.3.80"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "fbe734895e869dc429d78c4b433f8d17d95f8d05317440b4fad5ab2d33e596dc"
|
||||
dependencies = [
|
||||
"js-sys",
|
||||
"wasm-bindgen",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "webpki-roots"
|
||||
version = "0.26.11"
|
||||
|
||||
10
Cargo.toml
10
Cargo.toml
@@ -3,7 +3,6 @@ members = [
|
||||
"crates/engine-cpu",
|
||||
"crates/engine-gpu-cuda",
|
||||
"crates/engine-gpu-opencl",
|
||||
"crates/engine-montgomery",
|
||||
"crates/metrics",
|
||||
"crates/miner-cli",
|
||||
"crates/miner-service",
|
||||
@@ -17,7 +16,7 @@ resolver = "2"
|
||||
edition = "2021"
|
||||
authors = ["Quantus Network"]
|
||||
description = "Quantus External Miner Workspace"
|
||||
version = "0.2.1"
|
||||
version = "1.0.0"
|
||||
|
||||
[workspace.dependencies]
|
||||
anyhow = "1"
|
||||
@@ -31,17 +30,16 @@ 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 = "1.0.1", default-features = false }
|
||||
qpow-math = { git = "https://github.com/Quantus-Network/chain.git", package = "qpow-math", default-features = false }
|
||||
quantus-miner-api = { version = "0.0.3", default-features = false }
|
||||
rand = { version = "0.8.5", default-features = false }
|
||||
serde = { version = "1.0", features = ["derive"] }
|
||||
serde_json = { version = "1.0.132", default-features = false }
|
||||
sha2 = { version = "0.10", default-features = false }
|
||||
sha3 = { version = "0.10", default-features = false }
|
||||
thiserror = "1"
|
||||
tokio = { version = "1.36", features = ["full"] }
|
||||
warp = "0.3"
|
||||
# Use local pow-core crate instead of the external qpow-math; keep API compatibility at the engine layer.
|
||||
# qpow-math intentionally omitted here.
|
||||
# pow-core now delegates to qpow-math for core functionality while providing optimized mining scaffolding.
|
||||
|
||||
# GPU backend dependencies are declared in individual engine crates (no optional deps in workspace)
|
||||
|
||||
|
||||
@@ -16,7 +16,7 @@ default = ["baseline"]
|
||||
# Map engine features to pow-core features so engine-cpu consumers can toggle them.
|
||||
baseline = ["pow-core/baseline"]
|
||||
montgomery = ["pow-core/montgomery"]
|
||||
simd-sha3 = ["pow-core/simd-sha3"]
|
||||
simd-poseidon2 = ["pow-core/simd-poseidon2"]
|
||||
|
||||
[dependencies]
|
||||
pow-core = { path = "../pow-core" }
|
||||
@@ -27,3 +27,9 @@ anyhow = { workspace = true }
|
||||
|
||||
[dev-dependencies]
|
||||
hex = { workspace = true }
|
||||
criterion = "0.5"
|
||||
rand = { workspace = true, features = ["std", "std_rng"] }
|
||||
|
||||
[[bench]]
|
||||
name = "cpu_engine_bench"
|
||||
harness = false
|
||||
|
||||
56
crates/engine-cpu/benches/cpu_engine_bench.rs
Normal file
56
crates/engine-cpu/benches/cpu_engine_bench.rs
Normal file
@@ -0,0 +1,56 @@
|
||||
use criterion::{black_box, criterion_group, criterion_main, Criterion};
|
||||
use engine_cpu::{FastCpuEngine, MinerEngine, Range};
|
||||
use pow_core::{hash_from_nonce, JobContext};
|
||||
use primitive_types::U512;
|
||||
use rand::RngCore;
|
||||
use std::sync::atomic::AtomicBool;
|
||||
|
||||
fn bench_cpu_fast_engine(c: &mut Criterion) {
|
||||
// Create the engine
|
||||
let engine = FastCpuEngine::new();
|
||||
let cancel_flag = AtomicBool::new(false);
|
||||
|
||||
let large_range = Range {
|
||||
start: U512::from(1000u64),
|
||||
end: U512::from(101000u64), // 100,000 nonces
|
||||
};
|
||||
|
||||
c.bench_function("cpu_fast_large_range", |b| {
|
||||
b.iter(|| {
|
||||
let mut header = [0u8; 32];
|
||||
rand::thread_rng().fill_bytes(&mut header);
|
||||
let difficulty = U512::from(1000u64);
|
||||
let ctx = JobContext::new(header, difficulty);
|
||||
|
||||
let result = engine.search_range(
|
||||
black_box(&ctx),
|
||||
black_box(large_range.clone()),
|
||||
black_box(&cancel_flag),
|
||||
);
|
||||
black_box(result)
|
||||
})
|
||||
});
|
||||
}
|
||||
|
||||
fn bench_hash_from_nonce(c: &mut Criterion) {
|
||||
// Create a test job context
|
||||
let mut header = [0u8; 32];
|
||||
rand::thread_rng().fill_bytes(&mut header);
|
||||
let difficulty = U512::from(1000u64);
|
||||
let ctx = JobContext::new(header, difficulty);
|
||||
|
||||
// Create some test nonce values
|
||||
let test_nonce_values: Vec<U512> = (0..100).map(|i| U512::from(1000u64 + i)).collect();
|
||||
|
||||
c.bench_function("hash_from_nonce_single", |b| {
|
||||
let mut i = 0;
|
||||
b.iter(|| {
|
||||
let nonce = test_nonce_values[i % test_nonce_values.len()];
|
||||
i += 1;
|
||||
let hash = hash_from_nonce(black_box(&ctx), black_box(nonce));
|
||||
black_box(hash)
|
||||
})
|
||||
});
|
||||
}
|
||||
criterion_group!(benches, bench_cpu_fast_engine, bench_hash_from_nonce);
|
||||
criterion_main!(benches);
|
||||
@@ -14,7 +14,7 @@
|
||||
|
||||
use core::cmp::Ordering;
|
||||
|
||||
use pow_core::{distance_for_nonce, is_valid_distance, JobContext};
|
||||
use pow_core::{is_valid_nonce, JobContext};
|
||||
use primitive_types::U512;
|
||||
use std::sync::atomic::{AtomicBool, AtomicU64, Ordering as AtomicOrdering};
|
||||
use std::time::Duration;
|
||||
@@ -31,7 +31,7 @@ pub struct Range {
|
||||
pub struct Candidate {
|
||||
pub nonce: U512,
|
||||
pub work: [u8; 64], // big-endian representation of nonce
|
||||
pub distance: U512, // achieved distance for this nonce
|
||||
pub hash: U512, // output hash for this nonce
|
||||
}
|
||||
|
||||
/// Origin of a found candidate.
|
||||
@@ -78,8 +78,8 @@ pub trait MinerEngine: Send + Sync {
|
||||
/// Human-readable engine name (for logs/metrics).
|
||||
fn name(&self) -> &'static str;
|
||||
|
||||
/// Prepare a precomputed context for a job (header + threshold).
|
||||
fn prepare_context(&self, header_hash: [u8; 32], threshold: U512) -> JobContext;
|
||||
/// Prepare a precomputed context for a job (header + difficulty).
|
||||
fn prepare_context(&self, header_hash: [u8; 32], difficulty: U512) -> JobContext;
|
||||
|
||||
/// Search an inclusive nonce range with cancellation support.
|
||||
///
|
||||
@@ -109,8 +109,8 @@ impl MinerEngine for BaselineCpuEngine {
|
||||
"cpu-baseline"
|
||||
}
|
||||
|
||||
fn prepare_context(&self, header_hash: [u8; 32], threshold: U512) -> JobContext {
|
||||
JobContext::new(header_hash, threshold)
|
||||
fn prepare_context(&self, header_hash: [u8; 32], difficulty: U512) -> JobContext {
|
||||
JobContext::new(header_hash, difficulty)
|
||||
}
|
||||
|
||||
fn search_range(&self, ctx: &JobContext, range: Range, cancel: &AtomicBool) -> EngineStatus {
|
||||
@@ -128,17 +128,18 @@ impl MinerEngine for BaselineCpuEngine {
|
||||
return EngineStatus::Cancelled { hash_count };
|
||||
}
|
||||
|
||||
// Compute distance for this nonce using the context.
|
||||
let distance = distance_for_nonce(ctx, current);
|
||||
// Compute hash for this nonce using Bitcoin-style double Poseidon2
|
||||
let (is_valid, hash) =
|
||||
is_valid_nonce(ctx.header, current.to_big_endian(), ctx.difficulty);
|
||||
hash_count = hash_count.saturating_add(1);
|
||||
|
||||
// Check if it's valid under threshold.
|
||||
if is_valid_distance(ctx, distance) {
|
||||
// Check if it meets difficulty target
|
||||
if is_valid {
|
||||
let work = current.to_big_endian();
|
||||
let candidate = Candidate {
|
||||
nonce: current,
|
||||
work,
|
||||
distance,
|
||||
hash,
|
||||
};
|
||||
return EngineStatus::Found {
|
||||
candidate,
|
||||
@@ -178,12 +179,12 @@ impl MinerEngine for FastCpuEngine {
|
||||
"cpu-fast"
|
||||
}
|
||||
|
||||
fn prepare_context(&self, header_hash: [u8; 32], threshold: U512) -> JobContext {
|
||||
JobContext::new(header_hash, threshold)
|
||||
fn prepare_context(&self, header_hash: [u8; 32], difficulty: U512) -> JobContext {
|
||||
JobContext::new(header_hash, difficulty)
|
||||
}
|
||||
|
||||
fn search_range(&self, ctx: &JobContext, range: Range, cancel: &AtomicBool) -> EngineStatus {
|
||||
use pow_core::{distance_from_y, init_worker_y0, is_valid_distance, step_mul};
|
||||
use pow_core::{hash_from_nonce, is_valid_hash, step_nonce};
|
||||
|
||||
// Ensure start <= end (inclusive range). If not, treat as exhausted.
|
||||
if range.start > range.end {
|
||||
@@ -191,7 +192,6 @@ impl MinerEngine for FastCpuEngine {
|
||||
}
|
||||
|
||||
let mut current = range.start;
|
||||
let mut y = init_worker_y0(ctx, current);
|
||||
let mut hash_count: u64 = 0;
|
||||
|
||||
loop {
|
||||
@@ -200,17 +200,17 @@ impl MinerEngine for FastCpuEngine {
|
||||
return EngineStatus::Cancelled { hash_count };
|
||||
}
|
||||
|
||||
// Compute distance from current accumulator
|
||||
let distance = distance_from_y(ctx, y);
|
||||
// Compute hash using Bitcoin-style double Poseidon2
|
||||
let hash = hash_from_nonce(ctx, current);
|
||||
hash_count = hash_count.saturating_add(1);
|
||||
|
||||
if is_valid_distance(ctx, distance) {
|
||||
if is_valid_hash(ctx, hash) {
|
||||
let work = current.to_big_endian();
|
||||
return EngineStatus::Found {
|
||||
candidate: Candidate {
|
||||
nonce: current,
|
||||
work,
|
||||
distance,
|
||||
hash,
|
||||
},
|
||||
hash_count,
|
||||
origin: FoundOrigin::Cpu,
|
||||
@@ -218,13 +218,14 @@ impl MinerEngine for FastCpuEngine {
|
||||
}
|
||||
|
||||
if current == range.end {
|
||||
break EngineStatus::Exhausted { hash_count };
|
||||
break;
|
||||
}
|
||||
|
||||
// Advance to next nonce: y <- y * m (mod n), current <- current + 1
|
||||
y = step_mul(ctx, y);
|
||||
current = current.saturating_add(U512::one());
|
||||
// Advance to next nonce
|
||||
current = step_nonce(current);
|
||||
}
|
||||
|
||||
EngineStatus::Exhausted { hash_count }
|
||||
}
|
||||
}
|
||||
|
||||
@@ -259,9 +260,9 @@ impl MinerEngine for ChainManipulatorEngine {
|
||||
"cpu-chain-manipulator"
|
||||
}
|
||||
|
||||
fn prepare_context(&self, header_hash: [u8; 32], threshold: U512) -> JobContext {
|
||||
fn prepare_context(&self, header_hash: [u8; 32], difficulty: U512) -> JobContext {
|
||||
// Per-block throttling: do NOT increment here. We increment on Found (i.e., when a block is solved).
|
||||
let ctx = JobContext::new(header_hash, threshold);
|
||||
let ctx = JobContext::new(header_hash, difficulty);
|
||||
// Debug: log current throttle state at job start
|
||||
log::debug!(
|
||||
target: "miner",
|
||||
@@ -274,16 +275,14 @@ impl MinerEngine for ChainManipulatorEngine {
|
||||
}
|
||||
|
||||
fn search_range(&self, ctx: &JobContext, range: Range, cancel: &AtomicBool) -> EngineStatus {
|
||||
use pow_core::{distance_from_y, init_worker_y0, is_valid_distance, step_mul};
|
||||
use pow_core::{hash_from_nonce, is_valid_hash, step_nonce};
|
||||
|
||||
if range.start > range.end {
|
||||
return EngineStatus::Exhausted { hash_count: 0 };
|
||||
}
|
||||
|
||||
// Fast incremental path (same as cpu-fast) to start with high hashrate:
|
||||
// y0 = m^(h + start_nonce) mod n, then y = y * m (mod n) each step.
|
||||
// Bitcoin-style hashing path
|
||||
let mut current = range.start;
|
||||
let mut y = init_worker_y0(ctx, current);
|
||||
let mut hash_count: u64 = 0;
|
||||
let mut batch_counter: u64 = 0;
|
||||
|
||||
@@ -302,8 +301,8 @@ impl MinerEngine for ChainManipulatorEngine {
|
||||
return EngineStatus::Cancelled { hash_count };
|
||||
}
|
||||
|
||||
// Compute distance from current accumulator (incremental path).
|
||||
let distance = distance_from_y(ctx, y);
|
||||
// Compute hash using Bitcoin-style double Poseidon2
|
||||
let hash = hash_from_nonce(ctx, current);
|
||||
hash_count = hash_count.saturating_add(1);
|
||||
batch_counter = batch_counter.saturating_add(1);
|
||||
|
||||
@@ -311,7 +310,7 @@ impl MinerEngine for ChainManipulatorEngine {
|
||||
#[allow(unused_variables)]
|
||||
let _dbg_batch = batch_counter;
|
||||
|
||||
if is_valid_distance(ctx, distance) {
|
||||
if is_valid_hash(ctx, hash) {
|
||||
let work = current.to_big_endian();
|
||||
// Increment solved-block counter so the NEXT block throttles more.
|
||||
let _new_idx = self.job_index.fetch_add(1, AtomicOrdering::Relaxed) + 1;
|
||||
@@ -333,7 +332,7 @@ impl MinerEngine for ChainManipulatorEngine {
|
||||
candidate: Candidate {
|
||||
nonce: current,
|
||||
work,
|
||||
distance,
|
||||
hash,
|
||||
},
|
||||
hash_count,
|
||||
origin: FoundOrigin::Cpu,
|
||||
@@ -348,9 +347,7 @@ impl MinerEngine for ChainManipulatorEngine {
|
||||
|
||||
// Advance
|
||||
if current < range.end {
|
||||
// y <- y * m (mod n); current <- current + 1
|
||||
y = step_mul(ctx, y);
|
||||
current = current.saturating_add(U512::one());
|
||||
current = step_nonce(current);
|
||||
} else {
|
||||
break EngineStatus::Exhausted { hash_count };
|
||||
}
|
||||
@@ -371,60 +368,16 @@ mod tests {
|
||||
|
||||
fn make_ctx() -> JobContext {
|
||||
let header = [1u8; 32];
|
||||
let threshold = U512::MAX; // permissive threshold for "found" parity test
|
||||
JobContext::new(header, threshold)
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn baseline_and_fast_engines_find_same_candidate_on_small_range() {
|
||||
let ctx = make_ctx();
|
||||
|
||||
let range = Range {
|
||||
start: U512::from(0u64),
|
||||
end: U512::from(100u64),
|
||||
};
|
||||
|
||||
let cancel = AtomicBool::new(false);
|
||||
|
||||
let baseline = BaselineCpuEngine::new();
|
||||
let fast = FastCpuEngine::new();
|
||||
|
||||
let b_status = baseline.search_range(&ctx, range.clone(), &cancel);
|
||||
let f_status = fast.search_range(&ctx, range.clone(), &cancel);
|
||||
|
||||
match (b_status, f_status) {
|
||||
(
|
||||
EngineStatus::Found {
|
||||
candidate: b_cand,
|
||||
hash_count: b_hashes,
|
||||
origin: _,
|
||||
},
|
||||
EngineStatus::Found {
|
||||
candidate: f_cand,
|
||||
hash_count: f_hashes,
|
||||
origin: _,
|
||||
},
|
||||
) => {
|
||||
assert_eq!(
|
||||
b_cand.nonce, f_cand.nonce,
|
||||
"engines disagreed on winning nonce"
|
||||
);
|
||||
assert_eq!(
|
||||
b_cand.distance, f_cand.distance,
|
||||
"engines disagreed on distance"
|
||||
);
|
||||
assert_eq!(b_hashes, f_hashes, "engines disagreed on hash_count");
|
||||
}
|
||||
(b, f) => panic!("expected Found/Found, got baseline={b:?}, fast={f:?}"),
|
||||
}
|
||||
let difficulty = U512::from(1u64); // easy difficulty for "found" parity test
|
||||
JobContext::new(header, difficulty)
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn engine_returns_exhausted_when_no_solution_in_range() {
|
||||
// Use a very strict threshold to make solutions effectively impossible in a tiny range.
|
||||
// Use a very hard difficulty to make solutions effectively impossible in a tiny range.
|
||||
let header = [2u8; 32];
|
||||
let threshold = U512::zero();
|
||||
let ctx = JobContext::new(header, threshold);
|
||||
let difficulty = U512::MAX;
|
||||
let ctx = JobContext::new(header, difficulty);
|
||||
|
||||
let range = Range {
|
||||
start: U512::from(1u64),
|
||||
|
||||
@@ -32,6 +32,6 @@ cust = { version = "0.3", optional = true }
|
||||
rustacuda = { version = "0.1", optional = true }
|
||||
|
||||
# Host-side hashing and big-int precompute for GPU constants
|
||||
sha3 = { workspace = true }
|
||||
qp-poseidon-core = { workspace = true }
|
||||
num-bigint = "0.4"
|
||||
num-traits = { workspace = true }
|
||||
|
||||
@@ -1,6 +1,6 @@
|
||||
# engine-gpu-cuda (CUDA backend) – G1 bring‑up
|
||||
|
||||
This crate provides the CUDA GPU backend for the Quantus External Miner. It currently implements “G1” bring‑up: the per‑nonce modular multiply loop runs on the GPU (512‑bit Montgomery CIOS), while SHA3‑512 and the threshold check run on the host CPU. This allows correctness and plumbing to be validated before we move SHA3 and early‑exit onto the device in G2.
|
||||
This crate provides the CUDA GPU backend for the Quantus External Miner. It currently implements "G1" bring‑up: the per‑nonce modular multiply loop runs on the GPU (512‑bit Montgomery CIOS), while Poseidon2‑512 and the threshold check run on the host CPU. This allows correctness and plumbing to be validated before we move Poseidon2 and early‑exit onto the device in G2.
|
||||
|
||||
The backend is feature‑gated. When built with `--features cuda`, the crate’s build script compiles the CUDA kernel and embeds device images into the binary (CUBIN preferred, PTX as fallback). At runtime the engine selects an embedded image and launches the kernel to produce normalized `y` values per iteration.
|
||||
|
||||
@@ -10,16 +10,16 @@ The backend is feature‑gated. When built with `--features cuda`, the crate’s
|
||||
|
||||
- G1 (current):
|
||||
- Device: 512‑bit Montgomery multiply (CIOS using 64×64→128 via `__umul64hi`), maintaining `ŷ` in Montgomery domain and converting to normal domain for output.
|
||||
- Host: SHA3‑512 and threshold compare (and orchestration).
|
||||
- Host: Poseidon2‑512 and threshold compare (and orchestration).
|
||||
- Correctness: parity against CPU small‑range tests.
|
||||
- Performance: primarily limited by PCIe copy‑back and host SHA3. See “Tuning” below.
|
||||
- Performance: primarily limited by PCIe copy‑back and host Poseidon2. See "Tuning" below.
|
||||
|
||||
- G2 (next):
|
||||
- Device: SHA3‑512 (Keccak‑f[1600], 24 rounds) optimized for 64‑byte input.
|
||||
- Device: Poseidon2‑512 optimized for 64‑byte input.
|
||||
- Device: threshold compare + global early‑exit flag (atomic) + tiny candidate write.
|
||||
- Device: move constants to `__constant__` memory.
|
||||
- Host: poll early‑exit; no large copy‑backs (only candidate or counters).
|
||||
- Result: removes PCIe and host‑SHA3 bottlenecks; enables real GPU‑bound throughput. Selection will be enabled via `MINER_CUDA_MODE=g2` once available.
|
||||
- Result: removes PCIe and host‑Poseidon2 bottlenecks; enables real GPU‑bound throughput. Selection will be enabled via `MINER_CUDA_MODE=g2` once available.
|
||||
|
||||
---
|
||||
|
||||
@@ -43,7 +43,7 @@ Notes:
|
||||
|
||||
## Runtime selection and embeds
|
||||
|
||||
At startup, the engine prefers the embedded CUBIN; if absent it falls back to the embedded PTX. You can override with `MINER_CUDA_IMAGE=cubin|ptx`. To attempt the G2 path (device SHA3 + early-exit), set `MINER_CUDA_MODE=g2`; if the G2 kernel isn’t embedded/available for the current device, the engine will fall back to G1 automatically. You’ll see logs like:
|
||||
At startup, the engine prefers the embedded CUBIN; if absent it falls back to the embedded PTX. You can override with `MINER_CUDA_IMAGE=cubin|ptx`. To attempt the G2 path (device Poseidon2 + early-exit), set `MINER_CUDA_MODE=g2`; if the G2 kernel isn't embedded/available for the current device, the engine will fall back to G1 automatically. You'll see logs like:
|
||||
- `CUDA: using CUBIN (embedded)`
|
||||
- `CUDA: using PTX source = embedded`
|
||||
- (If neither exists, the engine logs the absence and delegates to CPU fast engine.)
|
||||
@@ -58,26 +58,26 @@ When a job runs, the engine prints its launch configuration and per‑launch out
|
||||
|
||||
## Env knobs – runtime (G1)
|
||||
|
||||
These knobs affect GPU launch shape and how much work is returned to the host (and thus how much SHA3 the CPU must perform per launch).
|
||||
These knobs affect GPU launch shape and how much work is returned to the host (and thus how much Poseidon2 the CPU must perform per launch).
|
||||
|
||||
- `MINER_CUDA_BLOCK_DIM` (default `256`)
|
||||
- Threads per block (`blockDim.x`). Use a multiple of 32 (warp size). 256 is a good default.
|
||||
- `MINER_CUDA_THREADS`
|
||||
- Total threads (grid workload). Grid dimension is `grid_dim = ceil(threads / block_dim)`. Target at least “#SMs × 1–2 blocks” for decent occupancy (e.g., RTX 3060 has 28 SMs → 28 or 32 blocks).
|
||||
- `MINER_CUDA_ITERS`
|
||||
- Iterations per thread. Higher values produce larger output buffers and more host SHA3 work per launch.
|
||||
- Iterations per thread. Higher values produce larger output buffers and more host Poseidon2 work per launch.
|
||||
- `MINER_CUDA_IMAGE` = `cubin` | `ptx` (optional)
|
||||
- Overrides the embedded image choice (debugging/testing). Default is to prefer CUBIN.
|
||||
- `MINER_CUDA_HASH_THREADS` (optional)
|
||||
- Number of host SHA3 worker threads to use to consume GPU output. Defaults to available parallelism.
|
||||
- Number of host Poseidon2 worker threads to use to consume GPU output. Defaults to available parallelism.
|
||||
- `MINER_CUDA_PINNED` = `1|true` (optional)
|
||||
- Use pinned (page-locked) host buffers and asynchronous device-to-host copies for G1 copy-back to reduce PCIe latency.
|
||||
- `MINER_CUDA_MODE` = `g2` (optional)
|
||||
- Attempt G2 kernel (device SHA3-512 + threshold compare + early-exit). Falls back to G1 if the G2 kernel is not available for the current device image.
|
||||
- Attempt G2 kernel (device Poseidon2-512 + threshold compare + early-exit). Falls back to G1 if the G2 kernel is not available for the current device image.
|
||||
|
||||
How much data per launch?
|
||||
- y_out bytes = `threads × iters × 64`.
|
||||
- Keep this around 64–128 MB in G1 to avoid PCIe and host SHA3 dominating.
|
||||
- Keep this around 64–128 MB in G1 to avoid PCIe and host Poseidon2 dominating.
|
||||
|
||||
Example configs (RTX 3060, SM 86):
|
||||
- ~64 MB per launch:
|
||||
@@ -115,12 +115,12 @@ The build fails with clear messages if:
|
||||
|
||||
## Tuning guide (G1)
|
||||
|
||||
Goal in G1: balance kernel time (GPU) against copy-back time (PCIe) and host SHA3 time (CPU) to avoid starving the GPU or overwhelming the host. Practical steps:
|
||||
Goal in G1: balance kernel time (GPU) against copy-back time (PCIe) and host Poseidon2 time (CPU) to avoid starving the GPU or overwhelming the host. Practical steps:
|
||||
|
||||
1) Size the output buffer:
|
||||
- Start with 64–128 MB per launch: `bytes ≈ threads × iters × 64`.
|
||||
- Increase `threads` to raise occupancy (more blocks). Start with `block_dim=256`.
|
||||
- Increase `iters` only while host SHA3 still keeps up.
|
||||
- Increase `iters` only while host Poseidon2 still keeps up.
|
||||
|
||||
2) Watch timings:
|
||||
- The engine logs `kernel_ms` and `copy_ms`.
|
||||
@@ -159,7 +159,7 @@ Goal in G1: balance kernel time (GPU) against copy-back time (PCIe) and host SHA
|
||||
|
||||
## Roadmap to G2
|
||||
|
||||
- Device SHA3‑512 (Keccak‑f[1600], 24 rounds) tuned for 64B input.
|
||||
- Device Poseidon2‑512 tuned for 64B input.
|
||||
- On‑device threshold compare and early‑exit flag (atomic).
|
||||
- Host polling and tiny candidate copy‑back.
|
||||
- Constants in `__constant__` memory.
|
||||
@@ -185,9 +185,9 @@ Runtime:
|
||||
- `MINER_CUDA_THREADS` — total threads (increase for more blocks).
|
||||
- `MINER_CUDA_ITERS` — iterations per thread (controls y_out size).
|
||||
- `MINER_CUDA_IMAGE` = `cubin|ptx` — force embedded image selection (optional).
|
||||
- `MINER_CUDA_HASH_THREADS` — parallel host SHA3 workers (optional).
|
||||
- `MINER_CUDA_HASH_THREADS` — parallel host Poseidon2 workers (optional).
|
||||
- `MINER_CUDA_PINNED` = `1|true` — use pinned host buffers + async D2H copy (G1 optimization).
|
||||
- `MINER_CUDA_MODE` = `g2` — try device SHA3 + early-exit; falls back to G1 if G2 kernel isn’t available.
|
||||
- `MINER_CUDA_MODE` = `g2` — try device Poseidon2 + early-exit; falls back to G1 if G2 kernel isn't available.
|
||||
|
||||
Build-time:
|
||||
- `CUDA_ARCH` = `sm_86|sm_89|sm_120|…` — SM target for device images (normalized internally).
|
||||
@@ -209,4 +209,4 @@ Presets are provided under `examples/.env` and follow a “lower” (≈1× SMs
|
||||
- RTX A5000: `cuda-miner-a5000-lower.env`, `cuda-miner-a5000-upper.env`
|
||||
- RTX A6000: `cuda-miner-a6000-lower.env`, `cuda-miner-a6000-upper.env`
|
||||
|
||||
Each preset uses `MINER_CUDA_MODE=g2` (device SHA3 + early-exit) and `MINER_CUDA_BLOCK_DIM=256`, and sizes `MINER_CUDA_THREADS` as `blocks × 256`. Adjust `MINER_CUDA_ITERS` to tune kernel dwell time vs early-exit responsiveness. If a G2 kernel image isn’t embedded for your device, the engine falls back to G1 automatically.
|
||||
Each preset uses `MINER_CUDA_MODE=g2` (device Poseidon2 + early-exit) and `MINER_CUDA_BLOCK_DIM=256`, and sizes `MINER_CUDA_THREADS` as `blocks × 256`. Adjust `MINER_CUDA_ITERS` to tune kernel dwell time vs early-exit responsiveness. If a G2 kernel image isn't embedded for your device, the engine falls back to G1 automatically.
|
||||
File diff suppressed because it is too large
Load Diff
@@ -1,101 +1,101 @@
|
||||
#![deny(rust_2018_idioms)]
|
||||
#![forbid(unsafe_code)]
|
||||
// #![deny(rust_2018_idioms)]
|
||||
// #![forbid(unsafe_code)]
|
||||
|
||||
//! OpenCL-based GPU mining engine (placeholder)
|
||||
//!
|
||||
//! This crate is a scaffold for a future OpenCL backend that will implement the
|
||||
//! mining engine interface used by the service layer. It currently provides:
|
||||
//! - An `OpenClEngine` type with a constructor and basic helpers.
|
||||
//! - Documentation of the intended integration points.
|
||||
//!
|
||||
//! Planned responsibilities (non-exhaustive):
|
||||
//! - Accept a prepared `JobContext` (from `pow-core`) per job.
|
||||
//! - Partition nonce ranges into GPU work assignments.
|
||||
//! - Run an OpenCL kernel that performs, per nonce in the range:
|
||||
//! - y <- y * m (mod n) using Montgomery multiplication (in Montgomery domain)
|
||||
//! - nonce_element <- SHA3_512(y) in the normal domain
|
||||
//! - distance <- target XOR nonce_element
|
||||
//! - if distance <= threshold: report solution and signal early-cancel
|
||||
//! - Coordinate early-exit via device-global flags and host polling.
|
||||
//!
|
||||
//! Notes:
|
||||
//! - This crate deliberately does NOT implement the `MinerEngine` trait yet,
|
||||
//! because the engine trait currently lives in `engine-cpu`. Once the trait
|
||||
//! is promoted to a shared crate (or re-exported for engines), this crate
|
||||
//! will implement it and become selectable at runtime via the service config.
|
||||
//! - OpenCL bindings (e.g., via the `ocl` crate) and kernels will be added
|
||||
//! behind feature flags (e.g., `opencl`). For now, we only offer placeholders
|
||||
//! so the workspace compiles cleanly and the integration points are clear.
|
||||
// //! OpenCL-based GPU mining engine (placeholder)
|
||||
// //!
|
||||
// //! This crate is a scaffold for a future OpenCL backend that will implement the
|
||||
// //! mining engine interface used by the service layer. It currently provides:
|
||||
// //! - An `OpenClEngine` type with a constructor and basic helpers.
|
||||
// //! - Documentation of the intended integration points.
|
||||
// //!
|
||||
// //! Planned responsibilities (non-exhaustive):
|
||||
// //! - Accept a prepared `JobContext` (from `pow-core`) per job.
|
||||
// //! - Partition nonce ranges into GPU work assignments.
|
||||
// //! - Run an OpenCL kernel that performs, per nonce in the range:
|
||||
// //! - y <- y * m (mod n) using Montgomery multiplication (in Montgomery domain)
|
||||
// //! - nonce_element <- SHA3_512(y) in the normal domain
|
||||
// //! - distance <- target XOR nonce_element
|
||||
// //! - if distance <= threshold: report solution and signal early-cancel
|
||||
// //! - Coordinate early-exit via device-global flags and host polling.
|
||||
// //!
|
||||
// //! Notes:
|
||||
// //! - This crate deliberately does NOT implement the `MinerEngine` trait yet,
|
||||
// //! because the engine trait currently lives in `engine-cpu`. Once the trait
|
||||
// //! is promoted to a shared crate (or re-exported for engines), this crate
|
||||
// //! will implement it and become selectable at runtime via the service config.
|
||||
// //! - OpenCL bindings (e.g., via the `ocl` crate) and kernels will be added
|
||||
// //! behind feature flags (e.g., `opencl`). For now, we only offer placeholders
|
||||
// //! so the workspace compiles cleanly and the integration points are clear.
|
||||
|
||||
use pow_core::JobContext;
|
||||
use primitive_types::U512;
|
||||
// use pow_core::JobContext;
|
||||
// use primitive_types::U512;
|
||||
|
||||
/// Placeholder type for the OpenCL engine.
|
||||
///
|
||||
/// When fully implemented, this engine will manage OpenCL platform/device
|
||||
/// discovery, context/queue creation, kernel compilation, memory transfers,
|
||||
/// and kernel launches. It will expose the same search-range semantics as
|
||||
/// the CPU engine(s) but backed by the GPU.
|
||||
#[derive(Default, Debug)]
|
||||
pub struct OpenClEngine {
|
||||
// Future fields (examples):
|
||||
// platform_id: usize,
|
||||
// device_id: usize,
|
||||
// context: ocl::Context,
|
||||
// queue: ocl::Queue,
|
||||
// program: ocl::Program,
|
||||
// kernel: ocl::Kernel,
|
||||
}
|
||||
// /// Placeholder type for the OpenCL engine.
|
||||
// ///
|
||||
// /// When fully implemented, this engine will manage OpenCL platform/device
|
||||
// /// discovery, context/queue creation, kernel compilation, memory transfers,
|
||||
// /// and kernel launches. It will expose the same search-range semantics as
|
||||
// /// the CPU engine(s) but backed by the GPU.
|
||||
// #[derive(Default, Debug)]
|
||||
// pub struct OpenClEngine {
|
||||
// // Future fields (examples):
|
||||
// // platform_id: usize,
|
||||
// // device_id: usize,
|
||||
// // context: ocl::Context,
|
||||
// // queue: ocl::Queue,
|
||||
// // program: ocl::Program,
|
||||
// // kernel: ocl::Kernel,
|
||||
// }
|
||||
|
||||
impl OpenClEngine {
|
||||
/// Construct a new OpenCL engine placeholder.
|
||||
///
|
||||
/// Future versions may accept configuration (e.g., platform/device index).
|
||||
pub fn new() -> Self {
|
||||
Self::default()
|
||||
}
|
||||
// impl OpenClEngine {
|
||||
// /// Construct a new OpenCL engine placeholder.
|
||||
// ///
|
||||
// /// Future versions may accept configuration (e.g., platform/device index).
|
||||
// pub fn new() -> Self {
|
||||
// Self::default()
|
||||
// }
|
||||
|
||||
/// Human-readable name for logs/metrics.
|
||||
pub fn name(&self) -> &'static str {
|
||||
"gpu-opencl (placeholder)"
|
||||
}
|
||||
// /// Human-readable name for logs/metrics.
|
||||
// pub fn name(&self) -> &'static str {
|
||||
// "gpu-opencl (placeholder)"
|
||||
// }
|
||||
|
||||
/// Prepare a precomputed job context for a given header and threshold.
|
||||
///
|
||||
/// This defers to `pow-core` to derive (m, n) and `target` from the header.
|
||||
/// In a full OpenCL implementation, this context will be uploaded to device
|
||||
/// constant buffers or passed as kernel arguments.
|
||||
pub fn prepare_context(&self, header_hash: [u8; 32], threshold: U512) -> JobContext {
|
||||
JobContext::new(header_hash, threshold)
|
||||
}
|
||||
// /// Prepare a precomputed job context for a given header and threshold.
|
||||
// ///
|
||||
// /// This defers to `pow-core` to derive (m, n) and `target` from the header.
|
||||
// /// In a full OpenCL implementation, this context will be uploaded to device
|
||||
// /// constant buffers or passed as kernel arguments.
|
||||
// pub fn prepare_context(&self, header_hash: [u8; 32], threshold: U512) -> JobContext {
|
||||
// JobContext::new(header_hash, threshold)
|
||||
// }
|
||||
|
||||
/// Returns whether this build has OpenCL support compiled in.
|
||||
///
|
||||
/// When actual OpenCL integration is added behind a feature flag, this will
|
||||
/// return true only if that feature is enabled.
|
||||
pub fn opencl_available(&self) -> bool {
|
||||
// Adjust once actual OpenCL integration is implemented behind a feature:
|
||||
// cfg!(feature = "opencl")
|
||||
false
|
||||
}
|
||||
}
|
||||
// /// Returns whether this build has OpenCL support compiled in.
|
||||
// ///
|
||||
// /// When actual OpenCL integration is added behind a feature flag, this will
|
||||
// /// return true only if that feature is enabled.
|
||||
// pub fn opencl_available(&self) -> bool {
|
||||
// // Adjust once actual OpenCL integration is implemented behind a feature:
|
||||
// // cfg!(feature = "opencl")
|
||||
// false
|
||||
// }
|
||||
// }
|
||||
|
||||
#[cfg(test)]
|
||||
mod tests {
|
||||
use super::*;
|
||||
use primitive_types::U512;
|
||||
// #[cfg(test)]
|
||||
// mod tests {
|
||||
// use super::*;
|
||||
// use primitive_types::U512;
|
||||
|
||||
#[test]
|
||||
fn placeholder_engine_basics() {
|
||||
let eng = OpenClEngine::new();
|
||||
assert_eq!(eng.name(), "gpu-opencl (placeholder)");
|
||||
// #[test]
|
||||
// fn placeholder_engine_basics() {
|
||||
// let eng = OpenClEngine::new();
|
||||
// assert_eq!(eng.name(), "gpu-opencl (placeholder)");
|
||||
|
||||
// Ensure context creation works and is deterministic in shape.
|
||||
let header = [1u8; 32];
|
||||
let threshold = U512::from(12345u64);
|
||||
let ctx = eng.prepare_context(header, threshold);
|
||||
// // Ensure context creation works and is deterministic in shape.
|
||||
// let header = [1u8; 32];
|
||||
// let threshold = U512::from(12345u64);
|
||||
// let ctx = eng.prepare_context(header, threshold);
|
||||
|
||||
assert_eq!(ctx.header, header);
|
||||
assert_eq!(ctx.threshold, threshold);
|
||||
}
|
||||
}
|
||||
// assert_eq!(ctx.header, header);
|
||||
// assert_eq!(ctx.threshold, threshold);
|
||||
// }
|
||||
// }
|
||||
|
||||
@@ -1,29 +0,0 @@
|
||||
[package]
|
||||
name = "engine-montgomery"
|
||||
version.workspace = true
|
||||
edition.workspace = true
|
||||
description = "Montgomery-optimized CPU mining engine for Quantus External Miner"
|
||||
|
||||
[lib]
|
||||
name = "engine_montgomery"
|
||||
path = "src/lib.rs"
|
||||
|
||||
[features]
|
||||
default = []
|
||||
metrics = ["dep:metrics"]
|
||||
|
||||
[dependencies]
|
||||
# Reuse the engine trait and shared types exposed by engine-cpu
|
||||
engine-cpu = { path = "../engine-cpu" }
|
||||
# Access job context and core PoW math interfaces
|
||||
pow-core = { path = "../pow-core" }
|
||||
|
||||
# Workspace-aligned deps
|
||||
primitive-types = { workspace = true }
|
||||
log = { workspace = true }
|
||||
sha3 = { workspace = true }
|
||||
num-traits = { workspace = true }
|
||||
metrics = { path = "../metrics", optional = true }
|
||||
|
||||
# Fixed-width big-integer backend for 512-bit Montgomery ops
|
||||
crypto-bigint = "0.5"
|
||||
@@ -1,210 +0,0 @@
|
||||
# engine-montgomery
|
||||
|
||||
Montgomery-optimized CPU mining engine for the Quantus External Miner.
|
||||
|
||||
This crate implements the same public `MinerEngine` trait as the baseline and fast CPU engines, while replacing the per-nonce modular multiplication with a fixed-width 8×64-bit Montgomery implementation. It is selectable via the CLI flag:
|
||||
|
||||
- `--engine cpu-montgomery`
|
||||
|
||||
The engine mirrors `cpu-fast` control flow and metrics emission so that apples-to-apples comparisons can be made between engines.
|
||||
|
||||
Highlights:
|
||||
- Fixed-width 512-bit arithmetic using 8×64-bit limbs.
|
||||
- Portable CIOS (Coarsely Integrated Operand Scanning) Montgomery multiply with `u128` intermediates.
|
||||
- Direct SHA3 over big-endian limbs to avoid intermediate big-integer conversions.
|
||||
- Per-job precompute cache (Montgomery params and `m_hat`) to reduce setup overhead.
|
||||
- Runtime backend selection for microarchitecture-optimized kernels (x86_64 BMI2-only and BMI2+ADX; aarch64 UMULH).
|
||||
- Metrics label for backend selection to aid dashboards and A/B analysis.
|
||||
|
||||
---
|
||||
|
||||
## Algorithm Overview
|
||||
|
||||
QPoW distance per nonce is computed as:
|
||||
|
||||
1) y update (group accumulation):
|
||||
- y0 = m^(h + start_nonce) mod n (one-time per worker).
|
||||
- y_{k+1} = (y_k * m) mod n (one multiply per nonce).
|
||||
|
||||
2) Distance = target XOR SHA3_512(y), where y is encoded as 64 big-endian bytes.
|
||||
|
||||
The engine focuses on optimizing step (1) with Montgomery multiplication and reducing overhead in step (2) by hashing directly over the final big-endian representation of y (no extra big-int intermediates).
|
||||
|
||||
---
|
||||
|
||||
## Montgomery Arithmetic (512-bit, 8×64 limbs)
|
||||
|
||||
We implement a 512-bit Montgomery field backed by 8 little-endian 64-bit limbs. For modulus `n` (odd composite, constrained by pow-core), we precompute:
|
||||
|
||||
- `n0_inv = -n^{-1} mod 2^64` (using Newton–Raphson).
|
||||
- `R = 2^(64*8) mod n` (implicit via representation).
|
||||
- `R^2 mod n` (computed once per job using the existing big-integer reference).
|
||||
- `to_mont(x) = x * R mod n = mont_mul(x, R^2)`.
|
||||
- `from_mont(x̂) = x̂ * 1 mod n = mont_mul(x̂, 1)`.
|
||||
|
||||
The core multiply is a portable CIOS Montgomery reduction:
|
||||
|
||||
```
|
||||
- acc <- 0
|
||||
- For each limb i in a:
|
||||
acc += a[i] * b
|
||||
m = (acc[0] * n0_inv) mod 2^64
|
||||
acc += m * n
|
||||
acc = acc >> 64 (shift down one limb)
|
||||
- If acc >= n: acc -= n
|
||||
- return acc
|
||||
```
|
||||
|
||||
Where `acc` is a 9-limb `u128` accumulator to simplify carries.
|
||||
|
||||
This yields a single multiply+reduce per nonce with a branchless inner loop (except final conditional subtract).
|
||||
|
||||
---
|
||||
|
||||
## Hashing Strategy
|
||||
|
||||
- We keep `y` in Montgomery domain during iteration to minimize transforms.
|
||||
- Before SHA3, we convert the residue `ŷ` back to the normal domain via `mont_mul(ŷ, 1)` and serialize as big-endian 64 bytes.
|
||||
- We reuse a single `Sha3_512` hasher per search call and `finalize_reset()` each iteration to reduce construction overhead.
|
||||
|
||||
This preserves consensus behavior while avoiding unnecessary big-int allocations or conversions.
|
||||
|
||||
---
|
||||
|
||||
## Backend Selection and Optimizations
|
||||
|
||||
At runtime, the engine selects a Montgomery multiply backend based on CPU and environment:
|
||||
|
||||
- Portable (default/fallback)
|
||||
- `mont_mul_portable`: pure-Rust, `u128`-based CIOS.
|
||||
- Available everywhere.
|
||||
|
||||
- x86_64 (runtime detected)
|
||||
- BMI2-only (`_mulx_u64`):
|
||||
- `mont_mul_bmi2`: uses BMI2 MULX to get 128-bit products efficiently.
|
||||
- Single carry chain (easier to validate; broadly available on newer CPUs).
|
||||
- BMI2+ADX:
|
||||
- `mont_mul_bmi2_adx`: implemented using MULX + ADCX/ADOX dual carry chains for higher ILP.
|
||||
|
||||
- aarch64
|
||||
- UMULH/ADCS:
|
||||
- `mont_mul_aarch64`: implemented using UMULH for high halves and ADCS-style accumulation via 64-bit ops to reduce dependency on `u128` where beneficial.
|
||||
- Default backend on Apple Silicon/macOS and Linux ARM64.
|
||||
|
||||
You can override backend selection for testing:
|
||||
|
||||
- `MINER_MONT_BACKEND=portable|bmi2|bmi2-adx|umulh`
|
||||
|
||||
The engine logs the selected backend (and exports it via metrics) at job start. Unsupported overrides safely fall back with a clear warning.
|
||||
|
||||
---
|
||||
|
||||
## Metrics
|
||||
|
||||
The engine emits the same per-job and per-thread metrics as other engines, but adds a backend info gauge:
|
||||
|
||||
- `miner_engine_backend{engine="cpu-montgomery", backend="<name>"} = 1`
|
||||
|
||||
This makes it easy to pivot in Grafana by backend.
|
||||
|
||||
All other metrics (hash rates, progress chunking cadence, counters) are identical to `cpu-fast` for apples-to-apples comparisons.
|
||||
|
||||
---
|
||||
|
||||
## Correctness and Tests
|
||||
|
||||
We keep the portable CIOS path as the ground-truth reference for optimized kernels, and we cross-check against pow-core’s BigUint-based implementations.
|
||||
|
||||
Property tests included:
|
||||
|
||||
- Portable Montgomery vs reference incremental multiply:
|
||||
- `from_mont(mul(to_mont(y), to_mont(m))) == step_mul(y)` across multiple steps.
|
||||
|
||||
- BMI2 vs Portable:
|
||||
- For randomized sequences, `bmi2` backend must match `portable` exactly at each step.
|
||||
- On non-x86_64 platforms, the test still runs but both backends fall back to `portable`.
|
||||
|
||||
- aarch64 UMULH vs Portable:
|
||||
- For randomized sequences, `aarch64-umulh` must match `portable` exactly at each step.
|
||||
- On non-aarch64 platforms, both tags fall back to `portable`.
|
||||
|
||||
- End-to-end parity:
|
||||
- `cpu-montgomery` vs `cpu-fast` on a small inclusive range (distance and winner parity; identical hash_count accounting).
|
||||
|
||||
We recommend running the property tests on machines with and without BMI2/ADX and on aarch64 to cover all optimized code paths.
|
||||
|
||||
---
|
||||
|
||||
## Safety
|
||||
|
||||
- The crate uses `#![deny(unsafe_code)]`.
|
||||
- `unsafe` is scoped only to tiny backend functions:
|
||||
- x86_64: `_mulx_u64` and inline asm for ADCX/ADOX dual carry chains.
|
||||
- aarch64: restricted intrinsics (e.g., UMULH) behind a small boundary.
|
||||
- All other code remains safe Rust.
|
||||
- The portable path is always available as a fallback for correctness/regression checks.
|
||||
|
||||
---
|
||||
|
||||
## Performance Notes
|
||||
|
||||
- Relative gains depend on how much time the miner spends in SHA3 vs modular multiply.
|
||||
- Direct-hash-from-residue + precompute caching already yields a measurable uplift over `cpu-fast`.
|
||||
- The BMI2-only path should improve throughput on supporting x86_64 hardware.
|
||||
- BMI2+ADX typically produces the highest gains on the multiply itself (often 1.5–2.0×), with end-to-end uplift bounded by SHA3 share per nonce.
|
||||
- aarch64 UMULH/ADCS brings similar relative gains on Apple Silicon and other ARM64 platforms.
|
||||
|
||||
To minimize orchestration overhead in the service:
|
||||
- Increase `--progress-chunk-ms` (e.g., 3000–5000) on both engines when comparing, to reduce update traffic and context switching.
|
||||
|
||||
---
|
||||
|
||||
## Runtime and CI Tips
|
||||
|
||||
- Selecting the engine:
|
||||
- `quantus-miner --engine cpu-montgomery`
|
||||
|
||||
- For A/B:
|
||||
- Keep workers and chunking identical across instances.
|
||||
- If testing backends explicitly:
|
||||
- `MINER_MONT_BACKEND=bmi2` or `bmi2-adx` on capable x86_64 hardware
|
||||
- `MINER_MONT_BACKEND=umulh` on aarch64
|
||||
- Check logs and metrics for the selected backend label.
|
||||
|
||||
- Observability:
|
||||
- Ensure metrics exporter is enabled (`--metrics-port ...`) for dashboards.
|
||||
- Filter or group by `engine="cpu-montgomery"` and backend metric to compare microarchitectural paths.
|
||||
|
||||
---
|
||||
|
||||
## Roadmap
|
||||
|
||||
- [x] Portable 8×64 CIOS (u128)
|
||||
- [x] Per-job precompute cache (`n0_inv`, `R^2 mod n`, `m_hat`)
|
||||
- [x] Direct SHA3 from normalized big-endian bytes
|
||||
- [x] Backend selection with log + metric
|
||||
- [x] x86_64 BMI2-only MULX kernel
|
||||
- [x] x86_64 BMI2+ADX (MULX + ADCX/ADOX) dual carry chain
|
||||
- [x] aarch64 UMULH/ADCS kernel (macOS/Linux ARM64)
|
||||
- [ ] Optional benchmark micro-harness (ns/op for mont_mul backends)
|
||||
- [ ] Extend tests with more randomized vectors and edge-case sweeps
|
||||
|
||||
---
|
||||
|
||||
## Design Rationale
|
||||
|
||||
- Keep the interface and metrics identical to `cpu-fast` so that any performance deltas reflect algorithmic/microarchitectural improvements rather than service overhead.
|
||||
- Keep a portable, well-reviewed core (CIOS) as a correctness reference.
|
||||
- Add microarchitecture-optimized kernels behind runtime dispatch and an env override for safe, controlled rollouts.
|
||||
- Log and export backend selection so A/B comparisons and regressions are easy to track.
|
||||
|
||||
---
|
||||
|
||||
## Contributing
|
||||
|
||||
- Changes to backends should include:
|
||||
- Property tests vs portable.
|
||||
- End-to-end parity checks vs `cpu-fast`.
|
||||
- A note in this README describing the optimization and any preconditions (e.g., required CPU features).
|
||||
- Keep unsafe code minimal, private, and well-commented.
|
||||
- Prefer small, focused PRs for each backend/optimization to simplify review and bisecting.
|
||||
@@ -1,974 +0,0 @@
|
||||
#![deny(rust_2018_idioms)]
|
||||
#![deny(unsafe_code)]
|
||||
|
||||
//! Montgomery-optimized CPU mining engine (scaffolding).
|
||||
//!
|
||||
//! Goals:
|
||||
//! - Mirror the cpu-fast engine behavior and metrics (hash counts, progress cadence).
|
||||
//! - Provide a drop-in engine selectable via `--engine cpu-montgomery`.
|
||||
//! - Introduce a crypto-bigint based 512-bit fixed-width backend scaffold for future
|
||||
//! Montgomery multiplication/reduction.
|
||||
//!
|
||||
//! Current state:
|
||||
//! - The search loop mirrors `engine-cpu`'s `FastCpuEngine` to ensure apples-for-apples
|
||||
//! metrics and correctness parity.
|
||||
//! - A lightweight Montgomery scaffolding is included (conversions and parameter
|
||||
//! container) using `crypto-bigint`, ready to be integrated into the hot path.
|
||||
//!
|
||||
//! Next steps (planned):
|
||||
//! - Replace the per-step `step_mul` with Montgomery domain multiplication:
|
||||
//! y_hat <- montgomery_mul(y_hat, m_hat, n, n') with y kept in Montgomery domain.
|
||||
//! - Precompute `R`, `R^2`, and `n'` once per job/thread and transform inputs.
|
||||
//! - Convert out of Montgomery before SHA3-512 distance computation.
|
||||
//!
|
||||
//! Important: We intentionally keep emissions (hash_count increments, control flow)
|
||||
//! identical to `cpu-fast` so metrics are directly comparable when pitting
|
||||
//! `--engine cpu-fast` vs `--engine cpu-montgomery`.
|
||||
|
||||
use core::cmp::Ordering;
|
||||
|
||||
use engine_cpu::EngineStatus;
|
||||
use engine_cpu::{EngineCandidate as Candidate, EngineRange as Range, MinerEngine};
|
||||
use pow_core::compat;
|
||||
use pow_core::{init_worker_y0, is_valid_distance, JobContext};
|
||||
use primitive_types::U512;
|
||||
|
||||
use std::collections::HashMap;
|
||||
use std::sync::atomic::{AtomicBool, Ordering as AtomicOrdering};
|
||||
use std::sync::{Arc, Mutex};
|
||||
|
||||
/// Montgomery engine for CPU.
|
||||
pub struct MontgomeryCpuEngine {
|
||||
#[allow(clippy::type_complexity)]
|
||||
cache: Mutex<HashMap<([u8; 64], [u8; 64]), Arc<mont_portable::MontCtx>>>,
|
||||
}
|
||||
|
||||
impl Default for MontgomeryCpuEngine {
|
||||
fn default() -> Self {
|
||||
Self {
|
||||
cache: Mutex::new(HashMap::new()),
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
impl MontgomeryCpuEngine {
|
||||
pub fn new() -> Self {
|
||||
Self::default()
|
||||
}
|
||||
}
|
||||
|
||||
impl MinerEngine for MontgomeryCpuEngine {
|
||||
fn name(&self) -> &'static str {
|
||||
"cpu-montgomery"
|
||||
}
|
||||
|
||||
fn prepare_context(&self, header_hash: [u8; 32], threshold: U512) -> JobContext {
|
||||
// Build the standard pow-core context (m, n, target, etc).
|
||||
// Montgomery parameters are computed per search call for now; we can
|
||||
// later thread them through a custom context wrapper if needed.
|
||||
JobContext::new(header_hash, threshold)
|
||||
}
|
||||
|
||||
fn search_range(&self, ctx: &JobContext, range: Range, cancel: &AtomicBool) -> EngineStatus {
|
||||
// NOTE: For initial integration we mirror the cpu-fast engine logic to ensure
|
||||
// metrics parity and correctness, while keeping the Montgomery scaffolding
|
||||
// ready for optimization in follow-ups.
|
||||
|
||||
if range.start > range.end {
|
||||
return EngineStatus::Exhausted { hash_count: 0 };
|
||||
}
|
||||
|
||||
// One-time init per thread: y0 = m^(h + start_nonce) mod n
|
||||
let mut current = range.start;
|
||||
let y0 = init_worker_y0(ctx, current);
|
||||
let mut hash_count: u64 = 0;
|
||||
|
||||
// Initialize or reuse per-job Montgomery params and residues from cache
|
||||
let mont = {
|
||||
let key = (ctx.m.to_big_endian(), ctx.n.to_big_endian());
|
||||
let mut guard = self.cache.lock().unwrap();
|
||||
if let Some(m) = guard.get(&key) {
|
||||
m.clone()
|
||||
} else {
|
||||
let m = Arc::new(mont_portable::MontCtx::from_ctx(ctx));
|
||||
guard.insert(key, m.clone());
|
||||
m
|
||||
}
|
||||
};
|
||||
let mut y_hat = mont.to_mont_u512(&y0);
|
||||
let m_hat = mont.m_hat;
|
||||
|
||||
loop {
|
||||
// Cancellation check (fast and frequent as in cpu-fast)
|
||||
if cancel.load(AtomicOrdering::Relaxed) {
|
||||
return EngineStatus::Cancelled { hash_count };
|
||||
}
|
||||
|
||||
// Compute distance from Montgomery accumulator: normalize then hash via pow-core
|
||||
let y_norm = mont.from_mont_u512(&y_hat);
|
||||
let distance = pow_core::distance_from_y(ctx, y_norm);
|
||||
hash_count = hash_count.saturating_add(1);
|
||||
|
||||
if is_valid_distance(ctx, distance) {
|
||||
let work = current.to_big_endian();
|
||||
return EngineStatus::Found {
|
||||
candidate: Candidate {
|
||||
nonce: current,
|
||||
work,
|
||||
distance,
|
||||
},
|
||||
hash_count,
|
||||
origin: engine_cpu::FoundOrigin::Cpu,
|
||||
};
|
||||
}
|
||||
|
||||
// Advance or finish
|
||||
match current.cmp(&range.end) {
|
||||
Ordering::Less => {
|
||||
// Incremental step in Montgomery domain: y_hat <- y_hat * m_hat (mod n)
|
||||
y_hat = mont.mul(&y_hat, &m_hat);
|
||||
current = current.saturating_add(U512::one());
|
||||
}
|
||||
_ => {
|
||||
break EngineStatus::Exhausted { hash_count };
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/// Montgomery scaffolding with crypto-bigint.
|
||||
/// This module provides conversions and parameter containers that we can use
|
||||
/// to wire up a fixed-width limb backend for 512-bit operations.
|
||||
mod mont_portable {
|
||||
use super::*;
|
||||
|
||||
// Montgomery context with portable CIOS 8x64 implementation (u128 intermediates).
|
||||
// Limbs are stored little-endian (limb 0 is least significant).
|
||||
type MulFn = fn(&[u64; 8], &[u64; 8], &[u64; 8], u64) -> [u64; 8];
|
||||
|
||||
#[derive(Clone)]
|
||||
pub struct MontCtx {
|
||||
n: [u64; 8],
|
||||
n0_inv: u64, // -n^{-1} mod 2^64
|
||||
r2: [u64; 8], // R^2 mod n
|
||||
pub m_hat: [u64; 8],
|
||||
mul_fn: MulFn,
|
||||
}
|
||||
|
||||
impl MontCtx {
|
||||
pub fn from_ctx(ctx: &JobContext) -> Self {
|
||||
let n = u512_to_le(ctx.n);
|
||||
let n0_inv = mont_n0_inv(n[0]);
|
||||
let r2_u512 = compat::mod_pow(&U512::from(2u32), &U512::from(1024u32), &ctx.n);
|
||||
let r2 = u512_to_le(r2_u512);
|
||||
let m = u512_to_le(ctx.m);
|
||||
let (mul_fn, backend) = select_backend();
|
||||
log::info!(target: "miner", "cpu-montgomery backend selected: {backend}");
|
||||
#[cfg(feature = "metrics")]
|
||||
{
|
||||
metrics::set_engine_backend("cpu-montgomery", backend);
|
||||
}
|
||||
let m_hat = mul_fn(&m, &r2, &n, n0_inv);
|
||||
MontCtx {
|
||||
n,
|
||||
n0_inv,
|
||||
r2,
|
||||
m_hat,
|
||||
mul_fn,
|
||||
}
|
||||
}
|
||||
|
||||
pub fn to_mont_u512(&self, x: &U512) -> [u64; 8] {
|
||||
let xl = u512_to_le(*x);
|
||||
(self.mul_fn)(&xl, &self.r2, &self.n, self.n0_inv)
|
||||
}
|
||||
|
||||
#[allow(clippy::wrong_self_convention)]
|
||||
pub fn from_mont_u512(&self, x_hat: &[u64; 8]) -> U512 {
|
||||
let one = {
|
||||
let mut o = [0u64; 8];
|
||||
o[0] = 1;
|
||||
o
|
||||
};
|
||||
let norm_le = (self.mul_fn)(x_hat, &one, &self.n, self.n0_inv);
|
||||
let norm_be = le_to_be_bytes(&norm_le);
|
||||
U512::from_big_endian(&norm_be)
|
||||
}
|
||||
|
||||
pub fn mul(&self, a_hat: &[u64; 8], b_hat: &[u64; 8]) -> [u64; 8] {
|
||||
(self.mul_fn)(a_hat, b_hat, &self.n, self.n0_inv)
|
||||
}
|
||||
|
||||
// Test-only helpers to enable forcing a specific backend and to access limb-level conversions.
|
||||
// These are useful for property tests and backend A/B validations.
|
||||
#[cfg(test)]
|
||||
pub fn from_ctx_with_backend_tag(ctx: &JobContext, tag: &str) -> Self {
|
||||
let n = u512_to_le(ctx.n);
|
||||
let n0_inv = mont_n0_inv(n[0]);
|
||||
let r2_u512 = compat::mod_pow(&U512::from(2u32), &U512::from(1024u32), &ctx.n);
|
||||
let r2 = u512_to_le(r2_u512);
|
||||
let m = u512_to_le(ctx.m);
|
||||
|
||||
// Choose mul_fn by tag; fall back to portable when not applicable or CPU features missing.
|
||||
let (mul_fn, _backend): (MulFn, &'static str) = match tag {
|
||||
"x86_64-bmi2-adx" | "bmi2-adx" => {
|
||||
#[cfg(target_arch = "x86_64")]
|
||||
{
|
||||
if std::is_x86_feature_detected!("bmi2")
|
||||
&& std::is_x86_feature_detected!("adx")
|
||||
{
|
||||
(mont_mul_bmi2_adx, "x86_64-bmi2-adx")
|
||||
} else if std::is_x86_feature_detected!("bmi2") {
|
||||
(mont_mul_bmi2, "x86_64-bmi2")
|
||||
} else {
|
||||
(mont_mul_portable, "portable")
|
||||
}
|
||||
}
|
||||
#[cfg(not(target_arch = "x86_64"))]
|
||||
{
|
||||
(mont_mul_portable, "portable")
|
||||
}
|
||||
}
|
||||
"x86_64-bmi2" | "bmi2" => {
|
||||
#[cfg(target_arch = "x86_64")]
|
||||
{
|
||||
if std::is_x86_feature_detected!("bmi2") {
|
||||
(mont_mul_bmi2, "x86_64-bmi2")
|
||||
} else {
|
||||
(mont_mul_portable, "portable")
|
||||
}
|
||||
}
|
||||
#[cfg(not(target_arch = "x86_64"))]
|
||||
{
|
||||
(mont_mul_portable, "portable")
|
||||
}
|
||||
}
|
||||
"aarch64-umulh" | "umulh" => {
|
||||
#[cfg(target_arch = "aarch64")]
|
||||
{
|
||||
(mont_mul_aarch64, "aarch64-umulh")
|
||||
}
|
||||
#[cfg(not(target_arch = "aarch64"))]
|
||||
{
|
||||
(mont_mul_portable, "portable")
|
||||
}
|
||||
}
|
||||
_ => (mont_mul_portable, "portable"),
|
||||
};
|
||||
|
||||
let m_hat = mul_fn(&m, &r2, &n, n0_inv);
|
||||
MontCtx {
|
||||
n,
|
||||
n0_inv,
|
||||
r2,
|
||||
m_hat,
|
||||
mul_fn,
|
||||
}
|
||||
}
|
||||
|
||||
#[cfg(test)]
|
||||
pub fn to_mont_le_limbs(&self, x: &U512) -> [u64; 8] {
|
||||
(self.mul_fn)(&u512_to_le(*x), &self.r2, &self.n, self.n0_inv)
|
||||
}
|
||||
|
||||
#[cfg(test)]
|
||||
#[allow(clippy::wrong_self_convention)]
|
||||
pub fn from_mont_le_limbs(&self, x_hat: &[u64; 8]) -> [u64; 8] {
|
||||
let one = {
|
||||
let mut o = [0u64; 8];
|
||||
o[0] = 1;
|
||||
o
|
||||
};
|
||||
(self.mul_fn)(x_hat, &one, &self.n, self.n0_inv)
|
||||
}
|
||||
}
|
||||
|
||||
#[inline]
|
||||
fn u512_to_le(x: U512) -> [u64; 8] {
|
||||
let be = x.to_big_endian();
|
||||
let mut limbs = [0u64; 8];
|
||||
// Split BE into 8 chunks, then reverse to get LE limb order (least-significant first).
|
||||
for i in 0..8 {
|
||||
let mut bytes = [0u8; 8];
|
||||
bytes.copy_from_slice(&be[i * 8..(i + 1) * 8]);
|
||||
limbs[i] = u64::from_be_bytes(bytes);
|
||||
}
|
||||
limbs.reverse();
|
||||
limbs
|
||||
}
|
||||
|
||||
#[inline]
|
||||
fn le_to_be_bytes(limbs: &[u64; 8]) -> [u8; 64] {
|
||||
let mut out = [0u8; 64];
|
||||
for i in 0..8 {
|
||||
let chunk = limbs[7 - i].to_be_bytes();
|
||||
out[i * 8..(i + 1) * 8].copy_from_slice(&chunk);
|
||||
}
|
||||
out
|
||||
}
|
||||
|
||||
// Compute n0_inv = -n[0]^{-1} mod 2^64 using Newton–Raphson (n[0] must be odd).
|
||||
#[inline]
|
||||
fn mont_n0_inv(n0: u64) -> u64 {
|
||||
// Compute inverse of n0 modulo 2^64
|
||||
let mut x = 1u64;
|
||||
// 6 iterations suffice for 64-bit modulus
|
||||
for _ in 0..6 {
|
||||
let t = x.wrapping_mul(n0);
|
||||
x = x.wrapping_mul(2u64.wrapping_sub(t));
|
||||
}
|
||||
x.wrapping_neg()
|
||||
}
|
||||
|
||||
// Portable CIOS Montgomery multiplication: returns (a * b * R^{-1}) mod n
|
||||
#[inline]
|
||||
fn mont_mul_portable(a: &[u64; 8], b: &[u64; 8], n: &[u64; 8], n0_inv: u64) -> [u64; 8] {
|
||||
const MASK: u128 = 0xFFFF_FFFF_FFFF_FFFFu128;
|
||||
let mut acc = [0u128; 9];
|
||||
|
||||
for &ai_u64 in a.iter().take(8) {
|
||||
// acc += ai * b
|
||||
let ai = ai_u64 as u128;
|
||||
let mut carry = 0u128;
|
||||
for j in 0..8 {
|
||||
let sum = acc[j] + ai * (b[j] as u128) + carry;
|
||||
acc[j] = sum & MASK;
|
||||
carry = sum >> 64;
|
||||
}
|
||||
acc[8] += carry;
|
||||
|
||||
// m = (acc[0] * n0_inv) mod 2^64
|
||||
let m = ((acc[0] as u64).wrapping_mul(n0_inv)) as u128;
|
||||
|
||||
// acc += m * n
|
||||
let mut carry2 = 0u128;
|
||||
for j in 0..8 {
|
||||
let sum = acc[j] + m * (n[j] as u128) + carry2;
|
||||
acc[j] = sum & MASK;
|
||||
carry2 = sum >> 64;
|
||||
}
|
||||
acc[8] += carry2;
|
||||
|
||||
// shift acc right by one limb
|
||||
for j in 0..8 {
|
||||
acc[j] = acc[j + 1];
|
||||
}
|
||||
acc[8] = 0;
|
||||
}
|
||||
|
||||
// Convert acc (little-endian limbs) to u64 array
|
||||
let mut res = [0u64; 8];
|
||||
for j in 0..8 {
|
||||
res[j] = acc[j] as u64;
|
||||
}
|
||||
|
||||
// Conditional subtract modulus if res >= n
|
||||
if ge_le(&res, n) {
|
||||
sub_le_in_place(&mut res, n);
|
||||
}
|
||||
|
||||
res
|
||||
}
|
||||
|
||||
#[inline]
|
||||
fn ge_le(a: &[u64; 8], b: &[u64; 8]) -> bool {
|
||||
for i in (0..8).rev() {
|
||||
if a[i] != b[i] {
|
||||
return a[i] > b[i];
|
||||
}
|
||||
}
|
||||
true
|
||||
}
|
||||
|
||||
#[inline]
|
||||
fn sub_le_in_place(a: &mut [u64; 8], b: &[u64; 8]) {
|
||||
let mut borrow: u128 = 0;
|
||||
for i in 0..8 {
|
||||
let ai = a[i] as u128;
|
||||
let bi = b[i] as u128;
|
||||
let tmp = (1u128 << 64) + ai - bi - borrow;
|
||||
a[i] = (tmp & 0xFFFF_FFFF_FFFF_FFFFu128) as u64;
|
||||
borrow = if tmp >> 64 == 0 { 1 } else { 0 };
|
||||
}
|
||||
}
|
||||
|
||||
#[inline]
|
||||
fn select_backend() -> (MulFn, &'static str) {
|
||||
// Optional override via env:
|
||||
// MINER_MONT_BACKEND=portable|bmi2|bmi2-adx
|
||||
if let Ok(val) = std::env::var("MINER_MONT_BACKEND") {
|
||||
let forced = val.to_ascii_lowercase();
|
||||
|
||||
#[cfg(target_arch = "x86_64")]
|
||||
{
|
||||
let bmi2 = std::is_x86_feature_detected!("bmi2");
|
||||
let adx = std::is_x86_feature_detected!("adx");
|
||||
match forced.as_str() {
|
||||
"portable" => {
|
||||
log::warn!(target: "miner", "cpu-montgomery backend override: forced portable");
|
||||
return (mont_mul_portable, "forced-portable");
|
||||
}
|
||||
"bmi2-adx" | "adx" => {
|
||||
if bmi2 && adx {
|
||||
log::warn!(target: "miner", "cpu-montgomery backend override: forced x86_64-bmi2-adx");
|
||||
return (mont_mul_bmi2_adx, "forced-x86_64-bmi2-adx");
|
||||
} else if bmi2 {
|
||||
log::warn!(target: "miner", "cpu-montgomery backend override requested bmi2-adx but ADX unavailable; falling back to x86_64-bmi2");
|
||||
return (mont_mul_bmi2, "x86_64-bmi2");
|
||||
} else {
|
||||
log::warn!(target: "miner", "cpu-montgomery backend override requested bmi2-adx but BMI2/ADX unavailable; falling back to x86_64-generic");
|
||||
return (mont_mul_portable, "x86_64-generic");
|
||||
}
|
||||
}
|
||||
"bmi2" => {
|
||||
if bmi2 {
|
||||
log::warn!(target: "miner", "cpu-montgomery backend override: forced x86_64-bmi2");
|
||||
return (mont_mul_bmi2, "forced-x86_64-bmi2");
|
||||
} else {
|
||||
log::warn!(target: "miner", "cpu-montgomery backend override requested bmi2 but BMI2 unavailable; falling back to x86_64-generic");
|
||||
return (mont_mul_portable, "x86_64-generic");
|
||||
}
|
||||
}
|
||||
other => {
|
||||
log::warn!(target: "miner", "cpu-montgomery backend override '{other}' is not recognized on x86_64; using auto-detect");
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
#[cfg(target_arch = "aarch64")]
|
||||
{
|
||||
match forced.as_str() {
|
||||
"portable" => {
|
||||
log::warn!(target: "miner", "cpu-montgomery backend override: forced portable");
|
||||
return (mont_mul_portable, "forced-portable");
|
||||
}
|
||||
// x86-only hints on aarch64 -> warn and ignore
|
||||
"bmi2" | "bmi2-adx" | "adx" => {
|
||||
log::warn!(target: "miner", "cpu-montgomery backend override '{}' not supported on aarch64; using auto-detect", forced);
|
||||
}
|
||||
other => {
|
||||
log::warn!(target: "miner", "cpu-montgomery backend override '{}' is not recognized on aarch64; using auto-detect", other);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
#[cfg(not(any(target_arch = "x86_64", target_arch = "aarch64")))]
|
||||
{
|
||||
match forced.as_str() {
|
||||
"portable" => {
|
||||
log::warn!(target: "miner", "cpu-montgomery backend override: forced portable");
|
||||
return (mont_mul_portable, "forced-portable");
|
||||
}
|
||||
other => {
|
||||
log::warn!(target: "miner", "cpu-montgomery backend override '{}' not supported on this arch; using auto-detect", other);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// Auto-detection (default) per architecture
|
||||
#[cfg(target_arch = "x86_64")]
|
||||
{
|
||||
let bmi2 = std::is_x86_feature_detected!("bmi2");
|
||||
|
||||
if bmi2 {
|
||||
(mont_mul_bmi2, "x86_64-bmi2")
|
||||
} else {
|
||||
(mont_mul_portable, "x86_64-generic")
|
||||
}
|
||||
}
|
||||
#[cfg(target_arch = "aarch64")]
|
||||
{
|
||||
(mont_mul_aarch64, "aarch64-umulh")
|
||||
}
|
||||
#[cfg(not(any(target_arch = "x86_64", target_arch = "aarch64")))]
|
||||
{
|
||||
(mont_mul_portable, "portable")
|
||||
}
|
||||
}
|
||||
|
||||
#[cfg(target_arch = "x86_64")]
|
||||
#[inline]
|
||||
#[allow(unsafe_code)]
|
||||
fn mont_mul_bmi2(a: &[u64; 8], b: &[u64; 8], n: &[u64; 8], n0_inv: u64) -> [u64; 8] {
|
||||
// BMI2-optimized CIOS using MULX. Single carry chain with u128 accumulation.
|
||||
// Limbs: little-endian (limb 0 = least significant).
|
||||
use core::arch::x86_64::_mulx_u64;
|
||||
const MASK: u128 = 0xFFFF_FFFF_FFFF_FFFFu128;
|
||||
|
||||
// 9-limb accumulator in u128 to simplify carries
|
||||
let mut acc = [0u128; 9];
|
||||
|
||||
for &ai in a.iter().take(8) {
|
||||
// acc += ai * b
|
||||
let mut carry: u128 = 0;
|
||||
for j in 0..8 {
|
||||
let mut hi: u64 = 0;
|
||||
// lo = (ai * b[j])_lo, hi = (ai * b[j])_hi
|
||||
let lo = unsafe { _mulx_u64(ai, b[j], &mut hi) };
|
||||
let sum = acc[j] + (lo as u128) + carry;
|
||||
acc[j] = sum & MASK;
|
||||
carry = (sum >> 64) + (hi as u128);
|
||||
}
|
||||
acc[8] += carry;
|
||||
|
||||
// m = (acc[0] * n0_inv) mod 2^64
|
||||
let m = (acc[0] as u64).wrapping_mul(n0_inv);
|
||||
|
||||
// acc += m * n
|
||||
let mut carry2: u128 = 0;
|
||||
for j in 0..8 {
|
||||
let mut hi2: u64 = 0;
|
||||
let lo2 = unsafe { _mulx_u64(m, n[j], &mut hi2) };
|
||||
let sum2 = acc[j] + (lo2 as u128) + carry2;
|
||||
acc[j] = sum2 & MASK;
|
||||
carry2 = (sum2 >> 64) + (hi2 as u128);
|
||||
}
|
||||
acc[8] += carry2;
|
||||
|
||||
// shift acc right by one limb (drop acc[0])
|
||||
for j in 0..8 {
|
||||
acc[j] = acc[j + 1];
|
||||
}
|
||||
acc[8] = 0;
|
||||
}
|
||||
|
||||
// Convert acc to u64 limbs (little-endian)
|
||||
let mut res = [0u64; 8];
|
||||
for j in 0..8 {
|
||||
res[j] = acc[j] as u64;
|
||||
}
|
||||
|
||||
// Conditional subtraction: if res >= n then res -= n
|
||||
if ge_le(&res, n) {
|
||||
sub_le_in_place(&mut res, n);
|
||||
}
|
||||
|
||||
res
|
||||
}
|
||||
|
||||
#[cfg(target_arch = "x86_64")]
|
||||
#[inline]
|
||||
#[allow(unsafe_code)]
|
||||
#[allow(unused_variables, unused_mut)]
|
||||
fn mont_mul_bmi2_adx(a: &[u64; 8], b: &[u64; 8], n: &[u64; 8], n0_inv: u64) -> [u64; 8] {
|
||||
// BMI2+ADX-optimized CIOS using MULX + dual carry chains (ADCX/ADOX).
|
||||
// We retain the same CIOS structure as the portable path:
|
||||
// - acc is a 9-limb accumulator in 64-bit limbs (little-endian)
|
||||
// - For each i:
|
||||
// acc += a[i] * b
|
||||
// m = (acc[0] * n0_inv) mod 2^64
|
||||
// acc += m * n
|
||||
// acc >>= 64 (drop acc[0])
|
||||
//
|
||||
// The inner adds use two independent carry chains:
|
||||
// - ADCX chain accumulates low halves into acc[j] (carry via CF)
|
||||
// - ADOX chain accumulates high halves into a separate running carry (OF)
|
||||
//
|
||||
// Note: We keep the logic in Rust around the asm! blocks for readability; the hot add paths
|
||||
// are emitted via inline assembly for ADX utilization.
|
||||
use core::arch::x86_64::_mulx_u64;
|
||||
use std::arch::asm;
|
||||
|
||||
let mut acc: [u64; 9] = [0; 9];
|
||||
|
||||
// helper: add (lo, hi) into acc[j] with dual carry chains (CF/OF)
|
||||
#[inline(always)]
|
||||
unsafe fn adx_accumulate(acc_j: &mut u64, lo: u64, hi: u64, of_carry: &mut u64) {
|
||||
// Clear CF and OF, then do:
|
||||
// acc_j += lo using ADCX (CF chain)
|
||||
// of_carry += hi using ADOX (OF chain)
|
||||
asm!(
|
||||
// Clear CF and OF carry chains
|
||||
"xor r8d, r8d",
|
||||
"adcx r8, r8",
|
||||
"adox r8, r8",
|
||||
// acc_j = acc_j + lo + CF
|
||||
"adcx {acc_j}, {lo}",
|
||||
// of_carry = of_carry + hi + OF
|
||||
"adox {ofc}, {hi}",
|
||||
acc_j = inout(reg) *acc_j,
|
||||
ofc = inout(reg) *of_carry,
|
||||
lo = in(reg) lo,
|
||||
hi = in(reg) hi,
|
||||
out("r8") _,
|
||||
options(nomem, nostack)
|
||||
);
|
||||
}
|
||||
|
||||
// helper: fold OF carry into next limb (acc[k] += of_carry)
|
||||
#[inline(always)]
|
||||
unsafe fn adx_fold_of(acc_k: &mut u64, of_carry: &mut u64) {
|
||||
// Fold the overflow-chain carry into acc_k using ADCX with cleared CF.
|
||||
asm!(
|
||||
"xor r8d, r8d",
|
||||
"adcx r8, r8",
|
||||
"adcx {acc_k}, {ofc}",
|
||||
acc_k = inout(reg) *acc_k,
|
||||
ofc = inout(reg) *of_carry,
|
||||
out("r8") _,
|
||||
options(nomem, nostack)
|
||||
);
|
||||
}
|
||||
|
||||
for &ai in a.iter().take(8) {
|
||||
// acc += ai * b
|
||||
let mut of_carry: u64 = 0;
|
||||
|
||||
// iterate j=0..7: accumulate ai*b[j] into acc[j].. with dual chains
|
||||
for j in 0..8 {
|
||||
let mut hi: u64 = 0;
|
||||
let lo: u64 = unsafe { _mulx_u64(ai, b[j], &mut hi) };
|
||||
|
||||
// acc[j] += lo (CF chain), of_carry += hi (OF chain)
|
||||
unsafe { adx_accumulate(&mut acc[j], lo, hi, &mut of_carry) };
|
||||
}
|
||||
// Propagate remaining OF carry into acc[8]
|
||||
unsafe { adx_fold_of(&mut acc[8], &mut of_carry) };
|
||||
|
||||
// m = (acc[0] * n0_inv) mod 2^64
|
||||
let m: u64 = (acc[0]).wrapping_mul(n0_inv);
|
||||
|
||||
// acc += m * n
|
||||
of_carry = 0;
|
||||
for j in 0..8 {
|
||||
let mut hi2: u64 = 0;
|
||||
let lo2: u64 = unsafe { _mulx_u64(m, n[j], &mut hi2) };
|
||||
|
||||
// acc[j] += lo2 (CF chain), of_carry += hi2 (OF chain)
|
||||
unsafe { adx_accumulate(&mut acc[j], lo2, hi2, &mut of_carry) };
|
||||
}
|
||||
// Propagate remaining OF carry into acc[8]
|
||||
unsafe { adx_fold_of(&mut acc[8], &mut of_carry) };
|
||||
|
||||
// shift acc right by one limb (drop acc[0])
|
||||
for j in 0..8 {
|
||||
acc[j] = acc[j + 1];
|
||||
}
|
||||
acc[8] = 0;
|
||||
}
|
||||
|
||||
// Conditional subtraction: if acc >= n then acc -= n
|
||||
let mut res = [0u64; 8];
|
||||
res.copy_from_slice(&acc[0..8]);
|
||||
|
||||
if ge_le(&res, n) {
|
||||
sub_le_in_place(&mut res, n);
|
||||
}
|
||||
|
||||
res
|
||||
}
|
||||
|
||||
#[cfg(target_arch = "aarch64")]
|
||||
#[inline]
|
||||
#[allow(unsafe_code)]
|
||||
fn mont_mul_aarch64(a: &[u64; 8], b: &[u64; 8], n: &[u64; 8], n0_inv: u64) -> [u64; 8] {
|
||||
const MASK: u128 = 0xFFFF_FFFF_FFFF_FFFFu128;
|
||||
let mut acc = [0u128; 9];
|
||||
for &ai in a.iter().take(8) {
|
||||
// acc += ai * b
|
||||
let mut carry: u128 = 0;
|
||||
for j in 0..8 {
|
||||
// low 64-bit product
|
||||
let lo = ai.wrapping_mul(b[j]);
|
||||
// high 64-bit product via UMULH intrinsic
|
||||
let hi = ((ai as u128) * (b[j] as u128)) >> 64;
|
||||
let sum = acc[j] + (lo as u128) + carry;
|
||||
acc[j] = sum & MASK;
|
||||
carry = (sum >> 64) + hi;
|
||||
}
|
||||
acc[8] += carry;
|
||||
|
||||
// m = (acc[0] * n0_inv) mod 2^64
|
||||
let m = (acc[0] as u64).wrapping_mul(n0_inv);
|
||||
|
||||
// acc += m * n
|
||||
let mut carry2: u128 = 0;
|
||||
for j in 0..8 {
|
||||
let lo2 = m.wrapping_mul(n[j]);
|
||||
let hi2 = ((m as u128) * (n[j] as u128)) >> 64;
|
||||
let sum2 = acc[j] + (lo2 as u128) + carry2;
|
||||
acc[j] = sum2 & MASK;
|
||||
carry2 = (sum2 >> 64) + hi2;
|
||||
}
|
||||
acc[8] += carry2;
|
||||
|
||||
// shift acc right by one limb
|
||||
for j in 0..8 {
|
||||
acc[j] = acc[j + 1];
|
||||
}
|
||||
acc[8] = 0;
|
||||
}
|
||||
|
||||
// Convert acc to result limbs
|
||||
let mut res = [0u64; 8];
|
||||
for j in 0..8 {
|
||||
res[j] = acc[j] as u64;
|
||||
}
|
||||
|
||||
// Conditional subtraction: if res >= n then res -= n
|
||||
if ge_le(&res, n) {
|
||||
sub_le_in_place(&mut res, n);
|
||||
}
|
||||
res
|
||||
}
|
||||
|
||||
#[cfg(test)]
|
||||
mod prop_tests {
|
||||
use super::*;
|
||||
use crate::MontgomeryCpuEngine;
|
||||
use engine_cpu::{EngineStatus, FastCpuEngine, MinerEngine};
|
||||
use pow_core::{init_worker_y0, step_mul, JobContext};
|
||||
use primitive_types::U512;
|
||||
use std::sync::atomic::AtomicBool;
|
||||
|
||||
fn u64x8_le_to_u512(le: &[u64; 8]) -> U512 {
|
||||
let be = le_to_be_bytes(le);
|
||||
U512::from_big_endian(&be)
|
||||
}
|
||||
|
||||
fn make_ctx_with_header_byte(byte: u8) -> JobContext {
|
||||
let mut header = [0u8; 32];
|
||||
header.fill(byte);
|
||||
let threshold = U512::MAX;
|
||||
JobContext::new(header, threshold)
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn montgomery_portable_mul_matches_step_mul() {
|
||||
// Validate that mont_mul agrees with pow_core's step_mul across many steps.
|
||||
let ctx = make_ctx_with_header_byte(0x5Au8);
|
||||
let mont = MontCtx::from_ctx_with_backend_tag(&ctx, "portable");
|
||||
|
||||
let start = U512::from(12345u64);
|
||||
let mut y_ref = init_worker_y0(&ctx, start);
|
||||
let mut y_hat = mont.to_mont_le_limbs(&y_ref);
|
||||
let m_hat = mont.m_hat;
|
||||
|
||||
// Walk 128 steps, comparing at each step
|
||||
for _ in 0..128 {
|
||||
// Reference path: y <- y * m (mod n) via pow_core BigUint
|
||||
y_ref = step_mul(&ctx, y_ref);
|
||||
|
||||
// Montgomery path: y_hat <- y_hat * m_hat (mod n), convert out
|
||||
y_hat = mont.mul(&y_hat, &m_hat);
|
||||
let y_hat_norm_le = mont.from_mont_le_limbs(&y_hat);
|
||||
let y_mont = u64x8_le_to_u512(&y_hat_norm_le);
|
||||
|
||||
assert_eq!(y_ref, y_mont, "montgomery mul mismatch vs step_mul");
|
||||
}
|
||||
}
|
||||
|
||||
#[cfg(target_arch = "x86_64")]
|
||||
#[test]
|
||||
fn montgomery_bmi2_equivalence_to_portable_when_available() {
|
||||
// On x86_64, ensure bmi2 path produces identical results as portable for the same (a,b).
|
||||
// On non-x86_64 this test still runs but both tags fall back to portable.
|
||||
let ctx = make_ctx_with_header_byte(0x3Cu8);
|
||||
let mont_port = MontCtx::from_ctx_with_backend_tag(&ctx, "portable");
|
||||
let mont_bmi2 = MontCtx::from_ctx_with_backend_tag(&ctx, "bmi2");
|
||||
|
||||
let start = U512::from(999u64);
|
||||
let mut y_ref = init_worker_y0(&ctx, start);
|
||||
|
||||
let mut y_hat_port = mont_port.to_mont_le_limbs(&y_ref);
|
||||
let mut y_hat_bmi2 = mont_bmi2.to_mont_le_limbs(&y_ref);
|
||||
|
||||
let m_hat_port = mont_port.m_hat;
|
||||
let m_hat_bmi2 = mont_bmi2.m_hat;
|
||||
|
||||
for _ in 0..64 {
|
||||
// advance reference so values change per-iteration
|
||||
y_ref = step_mul(&ctx, y_ref);
|
||||
|
||||
y_hat_port = mont_port.mul(&y_hat_port, &m_hat_port);
|
||||
y_hat_bmi2 = mont_bmi2.mul(&y_hat_bmi2, &m_hat_bmi2);
|
||||
|
||||
let y_port = u64x8_le_to_u512(&mont_port.from_mont_le_limbs(&y_hat_port));
|
||||
let y_bmi2 = u64x8_le_to_u512(&mont_bmi2.from_mont_le_limbs(&y_hat_bmi2));
|
||||
|
||||
assert_eq!(y_port, y_bmi2, "bmi2 path mismatch with portable");
|
||||
}
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn engine_end_to_end_matches_cpu_fast_on_small_range() {
|
||||
// End-to-end parity check against cpu-fast over a small inclusive range.
|
||||
let header = [0x11u8; 32];
|
||||
let threshold = U512::MAX;
|
||||
let ctx = JobContext::new(header, threshold);
|
||||
|
||||
let range = crate::Range {
|
||||
start: U512::from(0u64),
|
||||
end: U512::from(500u64),
|
||||
};
|
||||
let cancel = AtomicBool::new(false);
|
||||
|
||||
let mont = MontgomeryCpuEngine::new();
|
||||
let fast = FastCpuEngine::new();
|
||||
|
||||
let s_m = mont.search_range(&ctx, range.clone(), &cancel);
|
||||
let s_f = fast.search_range(&ctx, range.clone(), &cancel);
|
||||
|
||||
match (s_m, s_f) {
|
||||
(
|
||||
EngineStatus::Found {
|
||||
candidate: cm,
|
||||
hash_count: hm,
|
||||
origin: _,
|
||||
},
|
||||
EngineStatus::Found {
|
||||
candidate: cf,
|
||||
hash_count: hf,
|
||||
origin: _,
|
||||
},
|
||||
) => {
|
||||
assert_eq!(cm.nonce, cf.nonce, "nonce mismatch");
|
||||
assert_eq!(cm.distance, cf.distance, "distance mismatch");
|
||||
assert_eq!(hm, hf, "hash_count mismatch");
|
||||
}
|
||||
(
|
||||
EngineStatus::Exhausted { hash_count: hm },
|
||||
EngineStatus::Exhausted { hash_count: hf },
|
||||
) => {
|
||||
assert_eq!(hm, hf, "hash_count mismatch on Exhausted");
|
||||
}
|
||||
(m, f) => panic!("expected matching status, got mont={m:?}, fast={f:?}"),
|
||||
}
|
||||
}
|
||||
|
||||
#[cfg(target_arch = "aarch64")]
|
||||
#[test]
|
||||
fn montgomery_aarch64_equivalence_to_portable_when_available() {
|
||||
// On aarch64, ensure UMULH/ADCS path matches portable. On other arches this test
|
||||
// still runs but both tags fall back to portable.
|
||||
let ctx = make_ctx_with_header_byte(0x77u8);
|
||||
let mont_port = MontCtx::from_ctx_with_backend_tag(&ctx, "portable");
|
||||
let mont_arm = MontCtx::from_ctx_with_backend_tag(&ctx, "aarch64-umulh");
|
||||
|
||||
let start = U512::from(4242u64);
|
||||
let mut y_ref = init_worker_y0(&ctx, start);
|
||||
|
||||
let mut y_hat_port = mont_port.to_mont_le_limbs(&y_ref);
|
||||
let mut y_hat_arm = mont_arm.to_mont_le_limbs(&y_ref);
|
||||
|
||||
let m_hat_port = mont_port.m_hat;
|
||||
let m_hat_arm = mont_arm.m_hat;
|
||||
|
||||
for _ in 0..64 {
|
||||
// advance reference so values change per-iteration
|
||||
y_ref = step_mul(&ctx, y_ref);
|
||||
|
||||
y_hat_port = mont_port.mul(&y_hat_port, &m_hat_port);
|
||||
y_hat_arm = mont_arm.mul(&y_hat_arm, &m_hat_arm);
|
||||
|
||||
let y_port = u64x8_le_to_u512(&mont_port.from_mont_le_limbs(&y_hat_port));
|
||||
let y_arm = u64x8_le_to_u512(&mont_arm.from_mont_le_limbs(&y_hat_arm));
|
||||
|
||||
assert_eq!(y_port, y_arm, "aarch64 umulh path mismatch with portable");
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
#[cfg(test)]
|
||||
mod tests {
|
||||
use super::*;
|
||||
use std::sync::atomic::AtomicBool;
|
||||
|
||||
fn make_ctx() -> JobContext {
|
||||
let header = [1u8; 32];
|
||||
let threshold = U512::MAX; // permissive threshold for "found" parity test
|
||||
JobContext::new(header, threshold)
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn montgomery_engine_matches_fast_engine_on_small_range() {
|
||||
let ctx = make_ctx();
|
||||
|
||||
let range = Range {
|
||||
start: U512::from(0u64),
|
||||
end: U512::from(100u64),
|
||||
};
|
||||
|
||||
let cancel = AtomicBool::new(false);
|
||||
|
||||
let mont = MontgomeryCpuEngine::new();
|
||||
let fast = engine_cpu::FastCpuEngine::new();
|
||||
|
||||
let m_status = mont.search_range(&ctx, range.clone(), &cancel);
|
||||
let f_status = fast.search_range(&ctx, range.clone(), &cancel);
|
||||
|
||||
match (m_status, f_status) {
|
||||
(
|
||||
EngineStatus::Found {
|
||||
candidate: m_cand,
|
||||
hash_count: m_hashes,
|
||||
origin: _,
|
||||
},
|
||||
EngineStatus::Found {
|
||||
candidate: f_cand,
|
||||
hash_count: f_hashes,
|
||||
origin: _,
|
||||
},
|
||||
) => {
|
||||
assert_eq!(
|
||||
m_cand.nonce, f_cand.nonce,
|
||||
"engines disagreed on winning nonce"
|
||||
);
|
||||
assert_eq!(
|
||||
m_cand.distance, f_cand.distance,
|
||||
"engines disagreed on distance"
|
||||
);
|
||||
assert_eq!(m_hashes, f_hashes, "engines disagreed on hash_count");
|
||||
}
|
||||
(m, f) => panic!("expected Found/Found, got montgomery={m:?}, fast={f:?}"),
|
||||
}
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn engine_returns_exhausted_when_no_solution_in_range() {
|
||||
// Very strict threshold to avoid solutions in a tiny range.
|
||||
let header = [2u8; 32];
|
||||
let threshold = U512::zero();
|
||||
let ctx = JobContext::new(header, threshold);
|
||||
|
||||
let range = Range {
|
||||
start: U512::from(1u64),
|
||||
end: U512::from(1000u64),
|
||||
};
|
||||
|
||||
let cancel = AtomicBool::new(false);
|
||||
let eng = MontgomeryCpuEngine::new();
|
||||
|
||||
let status = eng.search_range(&ctx, range.clone(), &cancel);
|
||||
match status {
|
||||
EngineStatus::Exhausted { hash_count } => {
|
||||
// Inclusive range length = end - start + 1
|
||||
let expected = (range.end - range.start + U512::one()).as_u64();
|
||||
assert_eq!(hash_count, expected, "hash_count should equal range length");
|
||||
}
|
||||
other => panic!("expected Exhausted, got {other:?}"),
|
||||
}
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn engine_respects_immediate_cancellation() {
|
||||
let ctx = make_ctx();
|
||||
let range = Range {
|
||||
start: U512::from(0u64),
|
||||
end: U512::from(1_000_000u64),
|
||||
};
|
||||
let cancel = AtomicBool::new(true); // cancelled before starting
|
||||
let eng = MontgomeryCpuEngine::new();
|
||||
|
||||
let status = eng.search_range(&ctx, range, &cancel);
|
||||
match status {
|
||||
EngineStatus::Cancelled { hash_count } => {
|
||||
assert_eq!(hash_count, 0, "expected no work when cancelled immediately");
|
||||
}
|
||||
other => panic!("expected Cancelled, got {other:?}"),
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -40,7 +40,7 @@ struct Args {
|
||||
manip_throttle_cap: Option<u64>,
|
||||
|
||||
/// Mining engine to use (default: cpu-fast).
|
||||
/// Options: cpu-baseline, cpu-fast, cpu-chain-manipulator, cpu-montgomery, gpu-cuda, gpu-opencl
|
||||
/// Options: cpu-baseline, cpu-fast, cpu-chain-manipulator, gpu-cuda, gpu-opencl
|
||||
/// Note: GPU engines are currently unimplemented and will return a clear error at runtime.
|
||||
#[arg(long, env = "MINER_ENGINE", value_enum, default_value_t = EngineCli::CpuFast)]
|
||||
engine: EngineCli,
|
||||
@@ -96,8 +96,6 @@ enum EngineCli {
|
||||
CpuBaseline,
|
||||
/// Optimized CPU engine (incremental precompute + step_mul)
|
||||
CpuFast,
|
||||
/// Montgomery-optimized CPU engine (fixed-width 512-bit ops)
|
||||
CpuMontgomery,
|
||||
/// Throttling CPU engine that slows per block to help reduce difficulty
|
||||
CpuChainManipulator,
|
||||
/// CUDA GPU engine (unimplemented; selecting will return an error)
|
||||
@@ -111,7 +109,6 @@ impl From<EngineCli> for EngineSelection {
|
||||
match value {
|
||||
EngineCli::CpuBaseline => EngineSelection::CpuBaseline,
|
||||
EngineCli::CpuFast => EngineSelection::CpuFast,
|
||||
EngineCli::CpuMontgomery => EngineSelection::CpuMontgomery,
|
||||
EngineCli::CpuChainManipulator => EngineSelection::CpuChainManipulator,
|
||||
EngineCli::GpuCuda => EngineSelection::GpuCuda,
|
||||
EngineCli::GpuOpencl => EngineSelection::GpuOpenCl,
|
||||
|
||||
@@ -6,21 +6,18 @@ publish = false
|
||||
description = "Service layer: HTTP API, job orchestration, and engine abstraction for the Quantus External Miner"
|
||||
|
||||
[features]
|
||||
default = ["cpu", "metrics", "montgomery"]
|
||||
default = ["cpu", "metrics"]
|
||||
# Enable CPU engine by default.
|
||||
cpu = ["engine-cpu"]
|
||||
# Optional metrics/observability (Prometheus endpoint).
|
||||
metrics = [
|
||||
"dep:metrics",
|
||||
"engine-gpu-cuda?/metrics",
|
||||
"engine-montgomery?/metrics",
|
||||
"metrics/http-exporter",
|
||||
]
|
||||
# Optional GPU backends (off by default).
|
||||
cuda = ["dep:engine-gpu-cuda", "engine-gpu-cuda/cuda"]
|
||||
opencl = ["engine-gpu-opencl"]
|
||||
# Optional Montgomery-optimized CPU backend
|
||||
montgomery = ["engine-montgomery"]
|
||||
|
||||
[dependencies]
|
||||
# Workspace-shared deps
|
||||
@@ -42,7 +39,6 @@ quantus-miner-api = { workspace = true }
|
||||
# Local crates
|
||||
pow-core = { path = "../pow-core" }
|
||||
engine-cpu = { path = "../engine-cpu", optional = true }
|
||||
engine-montgomery = { path = "../engine-montgomery", optional = true }
|
||||
engine-gpu-cuda = { path = "../engine-gpu-cuda", optional = true, features = ["cuda"] }
|
||||
engine-gpu-opencl = { path = "../engine-gpu-opencl", optional = true }
|
||||
metrics = { path = "../metrics", optional = true }
|
||||
|
||||
@@ -3,7 +3,6 @@
|
||||
|
||||
use crossbeam_channel::{bounded, Receiver, Sender};
|
||||
use engine_cpu::{EngineCandidate, EngineRange, MinerEngine};
|
||||
use pow_core::compat;
|
||||
use primitive_types::U512;
|
||||
use quantus_miner_api::*;
|
||||
use std::collections::HashMap;
|
||||
@@ -45,7 +44,6 @@ pub struct ServiceConfig {
|
||||
pub enum EngineSelection {
|
||||
CpuBaseline,
|
||||
CpuFast,
|
||||
CpuMontgomery,
|
||||
CpuChainManipulator,
|
||||
GpuCuda,
|
||||
GpuOpenCl,
|
||||
@@ -72,7 +70,6 @@ impl fmt::Display for ServiceConfig {
|
||||
let engine = match self.engine {
|
||||
EngineSelection::CpuBaseline => "cpu-baseline",
|
||||
EngineSelection::CpuFast => "cpu-fast",
|
||||
EngineSelection::CpuMontgomery => "cpu-montgomery",
|
||||
EngineSelection::CpuChainManipulator => "cpu-chain-manipulator",
|
||||
EngineSelection::GpuCuda => "gpu-cuda",
|
||||
EngineSelection::GpuOpenCl => "gpu-opencl",
|
||||
@@ -281,14 +278,14 @@ pub enum JobStatus {
|
||||
pub struct MiningJobResult {
|
||||
pub nonce: U512,
|
||||
pub work: [u8; 64],
|
||||
pub distance: U512,
|
||||
pub hash: U512,
|
||||
}
|
||||
|
||||
/// Mining job data structure stored in the service.
|
||||
#[derive(Debug)]
|
||||
pub struct MiningJob {
|
||||
pub header_hash: [u8; 32],
|
||||
pub distance_threshold: U512,
|
||||
pub difficulty: U512,
|
||||
pub nonce_start: U512,
|
||||
pub nonce_end: U512,
|
||||
|
||||
@@ -313,7 +310,7 @@ impl Clone for MiningJob {
|
||||
fn clone(&self) -> Self {
|
||||
MiningJob {
|
||||
header_hash: self.header_hash,
|
||||
distance_threshold: self.distance_threshold,
|
||||
difficulty: self.difficulty,
|
||||
nonce_start: self.nonce_start,
|
||||
nonce_end: self.nonce_end,
|
||||
|
||||
@@ -349,13 +346,13 @@ pub struct ThreadResult {
|
||||
impl MiningJob {
|
||||
pub fn new(
|
||||
header_hash: [u8; 32],
|
||||
distance_threshold: U512,
|
||||
difficulty: U512,
|
||||
nonce_start: U512,
|
||||
nonce_end: U512,
|
||||
) -> Self {
|
||||
MiningJob {
|
||||
header_hash,
|
||||
distance_threshold,
|
||||
difficulty,
|
||||
nonce_start,
|
||||
nonce_end,
|
||||
status: JobStatus::Running,
|
||||
@@ -401,7 +398,7 @@ impl MiningJob {
|
||||
);
|
||||
|
||||
// Prepare shared job context once per job.
|
||||
let ctx = engine.prepare_context(self.header_hash, self.distance_threshold);
|
||||
let ctx = engine.prepare_context(self.header_hash, self.difficulty);
|
||||
|
||||
for (thread_id, (start, end)) in partitions.ranges.into_iter().enumerate() {
|
||||
let cancel_flag = self.cancel_flag.clone();
|
||||
@@ -537,20 +534,20 @@ impl MiningJob {
|
||||
let is_better = self
|
||||
.best_result
|
||||
.as_ref()
|
||||
.is_none_or(|current_best| result.distance < current_best.distance);
|
||||
.is_none_or(|current_best| result.hash < current_best.hash);
|
||||
|
||||
if is_better {
|
||||
log::debug!(target: "miner",
|
||||
"Found better result from thread {}: distance = {}, nonce = {}",
|
||||
thread_result.thread_id,
|
||||
result.distance,
|
||||
result.hash,
|
||||
result.nonce
|
||||
);
|
||||
self.best_result = Some(result.clone());
|
||||
self.cancel_flag.store(true, Ordering::Relaxed);
|
||||
// Result is now ready to be fetched via /result
|
||||
log::info!(target: "miner", "Result ready: engine={}, nonce={}, distance={}",
|
||||
self.engine_name, result.nonce, result.distance);
|
||||
self.engine_name, result.nonce, result.hash);
|
||||
#[cfg(feature = "metrics")]
|
||||
{
|
||||
// reuse existing http metric bucket for visibility until dedicated counters exist
|
||||
@@ -689,23 +686,14 @@ fn mine_range_with_engine(
|
||||
|
||||
match status {
|
||||
engine_cpu::EngineStatus::Found {
|
||||
candidate:
|
||||
EngineCandidate {
|
||||
nonce,
|
||||
work,
|
||||
distance,
|
||||
},
|
||||
candidate: EngineCandidate { nonce, work, hash },
|
||||
hash_count,
|
||||
origin,
|
||||
} => {
|
||||
// Send final result with found candidate and the hashes covered in this subrange
|
||||
let final_result = ThreadResult {
|
||||
thread_id,
|
||||
result: Some(MiningJobResult {
|
||||
nonce,
|
||||
work,
|
||||
distance,
|
||||
}),
|
||||
result: Some(MiningJobResult { nonce, work, hash }),
|
||||
hash_count,
|
||||
origin: Some(origin),
|
||||
completed: true,
|
||||
@@ -845,11 +833,11 @@ pub async fn handle_mine_request(
|
||||
.unwrap()
|
||||
.try_into()
|
||||
.expect("Validated hex string is 32 bytes");
|
||||
let distance_threshold = U512::from_dec_str(&request.distance_threshold).unwrap();
|
||||
let difficulty = U512::from_dec_str(&request.distance_threshold).unwrap();
|
||||
let nonce_start = U512::from_str_radix(&request.nonce_start, 16).unwrap();
|
||||
let nonce_end = U512::from_str_radix(&request.nonce_end, 16).unwrap();
|
||||
|
||||
let job = MiningJob::new(header_hash, distance_threshold, nonce_start, nonce_end);
|
||||
let job = MiningJob::new(header_hash, difficulty, nonce_start, nonce_end);
|
||||
|
||||
match state.add_job(request.job_id.clone(), job).await {
|
||||
Ok(_) => {
|
||||
@@ -941,16 +929,15 @@ pub async fn handle_result_request(
|
||||
// Inline re-verify using the exact nonce bytes we will return
|
||||
if let Some(result) = &job.best_result {
|
||||
let nonce_be = result.nonce.to_big_endian();
|
||||
let d2 = compat::get_nonce_distance(job.header_hash, nonce_be);
|
||||
let ok = d2 <= job.distance_threshold;
|
||||
let (ok, hash_result) = pow_core::is_valid_nonce(job.header_hash, nonce_be, job.difficulty);
|
||||
log::info!(
|
||||
target: "miner",
|
||||
"Serving result: job_id={}, engine={}, ok={}, host_distance={}, threshold={}",
|
||||
"Serving result: job_id={}, engine={}, ok={}, hash={}, difficulty={}",
|
||||
job_id,
|
||||
job.engine_name,
|
||||
ok,
|
||||
d2,
|
||||
job.distance_threshold
|
||||
hash_result,
|
||||
job.difficulty
|
||||
);
|
||||
#[cfg(feature = "metrics")]
|
||||
{
|
||||
@@ -1185,17 +1172,6 @@ pub async fn run(config: ServiceConfig) -> anyhow::Result<()> {
|
||||
let mut engine: Arc<dyn MinerEngine> = match config.engine {
|
||||
EngineSelection::CpuBaseline => Arc::new(engine_cpu::BaselineCpuEngine::new()),
|
||||
EngineSelection::CpuFast => Arc::new(engine_cpu::FastCpuEngine::new()),
|
||||
EngineSelection::CpuMontgomery => {
|
||||
#[cfg(feature = "montgomery")]
|
||||
{
|
||||
Arc::new(engine_montgomery::MontgomeryCpuEngine::new())
|
||||
}
|
||||
#[cfg(not(feature = "montgomery"))]
|
||||
{
|
||||
// Fallback if montgomery backend is not compiled in
|
||||
Arc::new(engine_cpu::FastCpuEngine::new())
|
||||
}
|
||||
}
|
||||
EngineSelection::CpuChainManipulator => {
|
||||
let mut eng = engine_cpu::ChainEngine::new();
|
||||
// Apply optional throttle parameters if provided.
|
||||
@@ -1437,13 +1413,13 @@ mod tests {
|
||||
let state = MiningService::new(2, engine, 2000);
|
||||
state.start_mining_loop().await;
|
||||
|
||||
// Impossible threshold with a nonce range that excludes 0
|
||||
// Impossible difficulty with a nonce range that excludes 0
|
||||
let header_hash = [1u8; 32];
|
||||
let distance_threshold = U512::zero();
|
||||
let difficulty = U512::MAX;
|
||||
let nonce_start = U512::from(1);
|
||||
let nonce_end = U512::from(100);
|
||||
|
||||
let job = MiningJob::new(header_hash, distance_threshold, nonce_start, nonce_end);
|
||||
let job = MiningJob::new(header_hash, difficulty, nonce_start, nonce_end);
|
||||
state.add_job("fail_job".to_string(), job).await.unwrap();
|
||||
|
||||
let mut finished_job = None;
|
||||
@@ -1535,13 +1511,13 @@ mod tests {
|
||||
let state = MiningService::new(2, engine, 2000);
|
||||
state.start_mining_loop().await;
|
||||
|
||||
// Easy threshold
|
||||
// Easy difficulty
|
||||
let header_hash = [1u8; 32];
|
||||
let distance_threshold = U512::MAX; // Easiest difficulty
|
||||
let difficulty = U512::from(1u64); // Easiest difficulty
|
||||
let nonce_start = U512::from(0);
|
||||
let nonce_end = U512::from(10000);
|
||||
|
||||
let job = MiningJob::new(header_hash, distance_threshold, nonce_start, nonce_end);
|
||||
let job = MiningJob::new(header_hash, difficulty, nonce_start, nonce_end);
|
||||
state.add_job("success_job".to_string(), job).await.unwrap();
|
||||
|
||||
let mut finished_job = None;
|
||||
@@ -1599,7 +1575,7 @@ mod tests {
|
||||
assert!(e.contains("mining_hash must be valid hex"));
|
||||
}
|
||||
|
||||
// 4) Bad distance_threshold decimal
|
||||
// 4) Bad difficulty decimal
|
||||
{
|
||||
let mut r = valid_req();
|
||||
r.distance_threshold = "not-a-decimal".to_string();
|
||||
@@ -1681,9 +1657,9 @@ mod tests {
|
||||
// 3) POST /mine valid -> 200 Accepted, duplicate -> 409
|
||||
let req = quantus_miner_api::MiningRequest {
|
||||
job_id: "job-http-1".to_string(),
|
||||
mining_hash: "11".repeat(32), // 64 hex chars
|
||||
distance_threshold: "0".to_string(), // strict, likely to fail later; OK for accept flow
|
||||
nonce_start: "00".repeat(64), // 128 hex chars
|
||||
mining_hash: "11".repeat(32), // 64 hex chars
|
||||
distance_threshold: "99999999999999".to_string(), // hard, likely to fail later; OK for accept flow
|
||||
nonce_start: "00".repeat(64), // 128 hex chars
|
||||
nonce_end: format!("{:0128x}", 1u8),
|
||||
};
|
||||
|
||||
@@ -1709,11 +1685,11 @@ mod tests {
|
||||
use crossbeam_channel::bounded;
|
||||
use std::sync::atomic::AtomicBool;
|
||||
|
||||
// Baseline engine, strict threshold to force Exhausted path for the sub-range
|
||||
// Baseline engine, hard difficulty to force Exhausted path for the sub-range
|
||||
let engine = engine_cpu::BaselineCpuEngine::new();
|
||||
let header = [3u8; 32];
|
||||
let threshold = U512::zero();
|
||||
let ctx = engine.prepare_context(header, threshold);
|
||||
let difficulty = U512::MAX;
|
||||
let ctx = engine.prepare_context(header, difficulty);
|
||||
|
||||
// Small range; chunking derives a large chunk size, so it will be a single chunk,
|
||||
// which still exercises the Exhausted -> progress update and final completion paths.
|
||||
|
||||
@@ -18,8 +18,7 @@ default = ["baseline", "std"]
|
||||
# std feature wires through to dependencies that have separate std/no_std builds.
|
||||
std = [
|
||||
"primitive-types/std",
|
||||
"sha2/std",
|
||||
"sha3/std",
|
||||
"qp-poseidon-core/std",
|
||||
]
|
||||
|
||||
# Baseline/reference implementation using BigUint-backed modular exponentiation.
|
||||
@@ -28,8 +27,8 @@ baseline = []
|
||||
# Enable accelerated modular arithmetic (e.g., Montgomery multiplication for 512-bit modulus).
|
||||
montgomery = []
|
||||
|
||||
# Optionally enable a faster SHA3/Keccak path (hook for SIMD/alt backends).
|
||||
simd-sha3 = []
|
||||
# Optionally enable a faster Poseidon2 path (hook for SIMD/alt backends).
|
||||
simd-poseidon2 = []
|
||||
|
||||
# Optional shared types/helpers for GPU engines (no device code here).
|
||||
gpu-interop = []
|
||||
@@ -39,8 +38,8 @@ log = { workspace = true }
|
||||
primitive-types = { workspace = true }
|
||||
num-bigint = { workspace = true }
|
||||
num-traits = { workspace = true }
|
||||
sha2 = { workspace = true }
|
||||
sha3 = { workspace = true }
|
||||
qp-poseidon-core = { workspace = true }
|
||||
qpow-math = { workspace = true }
|
||||
thiserror = { workspace = true }
|
||||
anyhow = { workspace = true }
|
||||
|
||||
|
||||
@@ -1,430 +1,192 @@
|
||||
#![cfg_attr(not(feature = "std"), no_std)]
|
||||
|
||||
// Local QPoW math core with a compatibility API mirroring the original `qpow-math` crate,
|
||||
// plus new scaffolding for an optimized path (precompute + incremental evaluation).
|
||||
//
|
||||
// This crate intentionally provides:
|
||||
// - `compat` API: Drop-in functions like `is_valid_nonce` and `get_nonce_distance`
|
||||
// - `JobContext`: Precomputed constants (m, n, target, threshold) for a given header
|
||||
// - Incremental helpers: `init_worker_y0`, `step_mul`, and `distance_from_y`
|
||||
// to enable replacing per-nonce exponentiation with a single modular multiplication.
|
||||
//
|
||||
// Notes:
|
||||
// - Current implementation mirrors the reference algorithm using BigUint-based modular arithmetic.
|
||||
// - Future work will gate accelerated paths (e.g., Montgomery, SIMD SHA3) behind features.
|
||||
|
||||
extern crate alloc;
|
||||
|
||||
use core::ops::BitXor;
|
||||
use primitive_types::U512;
|
||||
|
||||
#[cfg(feature = "std")]
|
||||
use log::{debug, error};
|
||||
pub use qpow_math::{get_nonce_hash, is_valid_nonce, mine_range};
|
||||
|
||||
pub mod compat {
|
||||
//! Compatibility layer that mirrors the original `qpow-math` crate API.
|
||||
|
||||
use super::*;
|
||||
|
||||
/// Check QPoW validity for a given `header` and `nonce` against `threshold`.
|
||||
///
|
||||
/// Returns a boolean indicating validity. If you also need the computed distance,
|
||||
/// use `is_valid_nonce_with_distance`.
|
||||
pub fn is_valid_nonce(header: [u8; 32], nonce: [u8; 64], threshold: U512) -> bool {
|
||||
let (ok, _) = is_valid_nonce_with_distance(header, nonce, threshold);
|
||||
ok
|
||||
}
|
||||
|
||||
/// Same as `is_valid_nonce`, but also returns the computed distance (U512).
|
||||
pub fn is_valid_nonce_with_distance(
|
||||
header: [u8; 32],
|
||||
nonce: [u8; 64],
|
||||
threshold: U512,
|
||||
) -> (bool, U512) {
|
||||
if nonce == [0u8; 64] {
|
||||
#[cfg(feature = "std")]
|
||||
error!(
|
||||
"is_valid_nonce should not be called with 0 nonce, but was for header: {header:?}"
|
||||
);
|
||||
return (false, U512::zero());
|
||||
}
|
||||
|
||||
let distance_achieved = get_nonce_distance(header, nonce);
|
||||
#[cfg(feature = "std")]
|
||||
debug!(target: "pow-core", "distance = {distance_achieved}..., threshold = {threshold}...");
|
||||
|
||||
(distance_achieved <= threshold, distance_achieved)
|
||||
}
|
||||
|
||||
/// Compute the QPoW distance for (header, nonce).
|
||||
///
|
||||
/// distance = target XOR H(m^(h + nonce) mod n)
|
||||
/// where (m, n) are derived deterministically from the header, and H is SHA3-512.
|
||||
pub fn get_nonce_distance(header: [u8; 32], nonce: [u8; 64]) -> U512 {
|
||||
super::get_nonce_distance_impl(header, nonce)
|
||||
}
|
||||
|
||||
/// Generate a pair (m, n) deterministically from the header.
|
||||
pub fn get_random_rsa(header: &[u8; 32]) -> (U512, U512) {
|
||||
super::get_random_rsa_impl(header)
|
||||
}
|
||||
|
||||
/// Check if two numbers are coprime using the Euclidean algorithm.
|
||||
pub fn is_coprime(a: &U512, b: &U512) -> bool {
|
||||
super::is_coprime_impl(a, b)
|
||||
}
|
||||
|
||||
/// Miller–Rabin primality test used by `get_random_rsa`.
|
||||
pub fn is_prime(n: &U512) -> bool {
|
||||
super::is_prime_impl(n)
|
||||
}
|
||||
|
||||
/// Apply the reference "hash-to-group" function then SHA3-512.
|
||||
pub fn hash_to_group_bigint_sha(h: &U512, m: &U512, n: &U512, solution: &U512) -> U512 {
|
||||
super::hash_to_group_bigint_sha_impl(h, m, n, solution)
|
||||
}
|
||||
|
||||
/// Reference hash-to-group function: computes m^(h + solution) mod n.
|
||||
pub fn hash_to_group_bigint(h: &U512, m: &U512, n: &U512, solution: &U512) -> U512 {
|
||||
super::hash_to_group_bigint_impl(h, m, n, solution)
|
||||
}
|
||||
|
||||
/// Reference modular exponentiation via BigUint.
|
||||
pub fn mod_pow(base: &U512, exponent: &U512, modulus: &U512) -> U512 {
|
||||
super::mod_pow_impl(base, exponent, modulus)
|
||||
}
|
||||
|
||||
/// SHA3-512 over the big-endian bytes of input U512.
|
||||
pub fn sha3_512(input: U512) -> U512 {
|
||||
super::sha3_512_impl(input)
|
||||
}
|
||||
}
|
||||
|
||||
/// Precomputed context for a single mining job (header + threshold).
|
||||
///
|
||||
/// This enables an optimized path:
|
||||
/// - Precompute (m, n) and target = H(m^(h + 0) mod n) once per job.
|
||||
/// - For each worker, compute `y0 = m^(h + start_nonce) mod n` once.
|
||||
/// - For each subsequent nonce, update y = (y * m) mod n (O(1) per step).
|
||||
/// - Distance at a step is `target XOR H(y)`.
|
||||
#[derive(Clone, Debug)]
|
||||
/// Job context for Bitcoin-style PoW mining with double Poseidon2 hashing
|
||||
#[derive(Debug, Clone)]
|
||||
pub struct JobContext {
|
||||
pub header: [u8; 32],
|
||||
pub header_int: U512,
|
||||
pub threshold: U512,
|
||||
|
||||
pub m: U512,
|
||||
pub n: U512,
|
||||
pub difficulty: U512,
|
||||
pub target: U512,
|
||||
}
|
||||
|
||||
impl JobContext {
|
||||
/// Build a new context by deriving (m, n) and target from the header.
|
||||
pub fn new(header: [u8; 32], threshold: U512) -> Self {
|
||||
let header_int = U512::from_big_endian(&header);
|
||||
let (m, n) = get_random_rsa_impl(&header);
|
||||
let target = hash_to_group_bigint_sha_impl(&header_int, &m, &n, &U512::zero());
|
||||
/// Build a new context from header and difficulty
|
||||
pub fn new(header: [u8; 32], difficulty: U512) -> Self {
|
||||
// In Bitcoin-style PoW, target = max_target / difficulty
|
||||
let max_target = U512::MAX;
|
||||
let target = max_target / difficulty;
|
||||
|
||||
JobContext {
|
||||
header,
|
||||
header_int,
|
||||
threshold,
|
||||
m,
|
||||
n,
|
||||
difficulty,
|
||||
target,
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/// Compute y0 = m^(h + start_nonce) mod n for a worker's starting nonce.
|
||||
///
|
||||
/// This is the one-time exponentiation cost per worker/thread. Subsequent steps can
|
||||
/// use `step_mul` to advance y with a single modular multiplication.
|
||||
pub fn init_worker_y0(ctx: &JobContext, start_nonce: U512) -> U512 {
|
||||
let sum = ctx.header_int.saturating_add(start_nonce);
|
||||
mod_pow_impl(&ctx.m, &sum, &ctx.n)
|
||||
/// Initialize a worker with starting nonce (no special initialization needed for Bitcoin-style)
|
||||
pub fn init_worker_nonce(start_nonce: U512) -> U512 {
|
||||
start_nonce
|
||||
}
|
||||
|
||||
/// Advance y by one nonce: y <- y * m (mod n).
|
||||
pub fn step_mul(ctx: &JobContext, y: U512) -> U512 {
|
||||
mod_mul_impl(&y, &ctx.m, &ctx.n)
|
||||
/// Advance nonce by one (simple increment for Bitcoin-style)
|
||||
pub fn step_nonce(nonce: U512) -> U512 {
|
||||
nonce.saturating_add(U512::from(1u64))
|
||||
}
|
||||
|
||||
/// Compute distance for the current y:
|
||||
/// distance = target XOR SHA3_512(y)
|
||||
pub fn distance_from_y(ctx: &JobContext, y: U512) -> U512 {
|
||||
let hashed = sha3_512_impl(y);
|
||||
ctx.target.bitxor(hashed)
|
||||
/// Compute hash for the current nonce using Bitcoin-style double Poseidon2
|
||||
pub fn hash_from_nonce(ctx: &JobContext, nonce: U512) -> U512 {
|
||||
let nonce_bytes = nonce.to_big_endian();
|
||||
qpow_math::get_nonce_hash(ctx.header, nonce_bytes)
|
||||
}
|
||||
|
||||
/// Convenience: compute distance for an arbitrary nonce using the context.
|
||||
pub fn distance_for_nonce(ctx: &JobContext, nonce: U512) -> U512 {
|
||||
let y = mod_pow_impl(&ctx.m, &ctx.header_int.saturating_add(nonce), &ctx.n);
|
||||
distance_from_y(ctx, y)
|
||||
/// Check if hash meets difficulty target
|
||||
pub fn is_valid_hash(ctx: &JobContext, hash: U512) -> bool {
|
||||
hash < ctx.target
|
||||
}
|
||||
|
||||
/// Convenience: check if a distance is valid under the context's threshold.
|
||||
pub fn is_valid_distance(ctx: &JobContext, distance: U512) -> bool {
|
||||
distance <= ctx.threshold
|
||||
}
|
||||
/// Mine a range of nonces starting from start_nonce
|
||||
pub fn mine_nonce_range(ctx: &JobContext, start_nonce: U512, steps: u64) -> Option<(U512, U512)> {
|
||||
let start_nonce_bytes = start_nonce.to_big_endian();
|
||||
|
||||
/// Reference distance computation used by the compat layer.
|
||||
fn get_nonce_distance_impl(header: [u8; 32], nonce: [u8; 64]) -> U512 {
|
||||
if nonce == [0u8; 64] {
|
||||
#[cfg(feature = "std")]
|
||||
debug!(target: "pow-core", "zero nonce");
|
||||
return U512::zero();
|
||||
if let Some((nonce_bytes, hash)) =
|
||||
mine_range(ctx.header, start_nonce_bytes, steps, ctx.difficulty)
|
||||
{
|
||||
let nonce = U512::from_big_endian(&nonce_bytes);
|
||||
Some((nonce, hash))
|
||||
} else {
|
||||
None
|
||||
}
|
||||
|
||||
let (m, n) = get_random_rsa_impl(&header);
|
||||
let header_int = U512::from_big_endian(&header);
|
||||
let nonce_int = U512::from_big_endian(&nonce);
|
||||
|
||||
let target = hash_to_group_bigint_sha_impl(&header_int, &m, &n, &U512::zero());
|
||||
let nonce_element = hash_to_group_bigint_sha_impl(&header_int, &m, &n, &nonce_int);
|
||||
|
||||
let distance = target.bitxor(nonce_element);
|
||||
#[cfg(feature = "std")]
|
||||
debug!(target: "pow-core", "distance = {distance}");
|
||||
distance
|
||||
}
|
||||
|
||||
/// Generates a pair of RSA-style numbers (m, n) deterministically from input header.
|
||||
///
|
||||
/// - m: 256-bit derived via SHA2-256(header)
|
||||
/// - n: 512-bit derived via SHA3-512(header), iteratively rehashed until valid:
|
||||
/// (odd, composite, coprime with m, and n > m)
|
||||
fn get_random_rsa_impl(header: &[u8; 32]) -> (U512, U512) {
|
||||
use sha2::{Digest, Sha256};
|
||||
use sha3::Sha3_512;
|
||||
|
||||
// m from SHA2-256
|
||||
let mut sha256 = Sha256::new();
|
||||
sha256.update(header);
|
||||
let m = U512::from_big_endian(sha256.finalize().as_slice());
|
||||
|
||||
// initial n from SHA3-512
|
||||
let mut sha3 = Sha3_512::new();
|
||||
sha3.update(header);
|
||||
let mut n = U512::from_big_endian(sha3.finalize().as_slice());
|
||||
|
||||
// Keep hashing until n satisfies constraints
|
||||
while n % 2u32 == U512::zero() || n <= m || !is_coprime_impl(&m, &n) || is_prime_impl(&n) {
|
||||
n = sha3_512_impl(n);
|
||||
}
|
||||
|
||||
(m, n)
|
||||
}
|
||||
|
||||
/// Check if two numbers are coprime using the Euclidean algorithm.
|
||||
fn is_coprime_impl(a: &U512, b: &U512) -> bool {
|
||||
let mut x = *a;
|
||||
let mut y = *b;
|
||||
|
||||
while y != U512::zero() {
|
||||
let tmp = y;
|
||||
y = x % y;
|
||||
x = tmp;
|
||||
}
|
||||
|
||||
x == U512::one()
|
||||
}
|
||||
|
||||
/// Hash-to-group then SHA3-512.
|
||||
///
|
||||
/// Note: The reference calls `hash_to_group_bigint` followed by an additional SHA3-512.
|
||||
fn hash_to_group_bigint_sha_impl(h: &U512, m: &U512, n: &U512, solution: &U512) -> U512 {
|
||||
let result = hash_to_group_bigint_impl(h, m, n, solution);
|
||||
sha3_512_impl(result)
|
||||
}
|
||||
|
||||
/// Reference hash-to-group big-integer function (no chunk splitting).
|
||||
/// Computes sum = h + solution; then y = m^sum mod n.
|
||||
fn hash_to_group_bigint_impl(h: &U512, m: &U512, n: &U512, solution: &U512) -> U512 {
|
||||
let sum = h.saturating_add(*solution);
|
||||
mod_pow_impl(m, &sum, n)
|
||||
}
|
||||
|
||||
/// Reference modular exponentiation using BigUint square-and-multiply.
|
||||
fn mod_pow_impl(base: &U512, exponent: &U512, modulus: &U512) -> U512 {
|
||||
if *modulus == U512::zero() {
|
||||
panic!("Modulus cannot be zero");
|
||||
}
|
||||
|
||||
use num_bigint::BigUint;
|
||||
use num_traits::{One, Zero};
|
||||
|
||||
// Convert inputs to BigUint
|
||||
let mut base = BigUint::from_bytes_be(&base.to_big_endian());
|
||||
let mut exp = BigUint::from_bytes_be(&exponent.to_big_endian());
|
||||
let modulus = BigUint::from_bytes_be(&modulus.to_big_endian());
|
||||
|
||||
// Initialize result as 1
|
||||
let mut result = BigUint::one();
|
||||
|
||||
// Square-and-multiply algorithm
|
||||
while !exp.is_zero() {
|
||||
if exp.bit(0) {
|
||||
result = (result * &base) % &modulus;
|
||||
}
|
||||
base = (&base * &base) % &modulus;
|
||||
exp >>= 1;
|
||||
}
|
||||
|
||||
U512::from_big_endian(&result.to_bytes_be())
|
||||
}
|
||||
|
||||
/// Reference modular multiplication using BigUint, i.e., (a * b) mod n.
|
||||
fn mod_mul_impl(a: &U512, b: &U512, modulus: &U512) -> U512 {
|
||||
use num_bigint::BigUint;
|
||||
|
||||
if *modulus == U512::zero() {
|
||||
panic!("Modulus cannot be zero");
|
||||
}
|
||||
|
||||
let a_bi = BigUint::from_bytes_be(&a.to_big_endian());
|
||||
let b_bi = BigUint::from_bytes_be(&b.to_big_endian());
|
||||
let n_bi = BigUint::from_bytes_be(&modulus.to_big_endian());
|
||||
|
||||
let prod = (a_bi * b_bi) % n_bi;
|
||||
U512::from_big_endian(&prod.to_bytes_be())
|
||||
}
|
||||
|
||||
/// Miller–Rabin primality test.
|
||||
///
|
||||
/// Deterministically selects k=32 bases hashed from `n` using SHA3-512 to
|
||||
/// bound false-positive probability to ~1/2^64 for composites.
|
||||
fn is_prime_impl(n: &U512) -> bool {
|
||||
use sha3::{Digest, Sha3_512};
|
||||
|
||||
if *n <= U512::one() {
|
||||
return false;
|
||||
}
|
||||
if *n == U512::from(2u32) || *n == U512::from(3u32) {
|
||||
return true;
|
||||
}
|
||||
if *n % U512::from(2u32) == U512::zero() {
|
||||
return false;
|
||||
}
|
||||
|
||||
// write n-1 as d * 2^r
|
||||
let mut d = *n - U512::one();
|
||||
let mut r = 0u32;
|
||||
while d % U512::from(2u32) == U512::zero() {
|
||||
d /= U512::from(2u32);
|
||||
r += 1;
|
||||
}
|
||||
|
||||
// Generate test bases deterministically from n using SHA3
|
||||
let mut bases = [U512::zero(); 32];
|
||||
let mut base_count = 0;
|
||||
let mut sha3 = Sha3_512::new();
|
||||
let mut counter = U512::zero();
|
||||
|
||||
while base_count < 32 {
|
||||
// Hash n concatenated with counter
|
||||
let mut bytes = [0u8; 128];
|
||||
let n_bytes = n.to_big_endian();
|
||||
let counter_bytes = counter.to_big_endian();
|
||||
|
||||
bytes[..64].copy_from_slice(&n_bytes);
|
||||
bytes[64..128].copy_from_slice(&counter_bytes);
|
||||
|
||||
sha3.update(bytes);
|
||||
|
||||
// Use the hash to generate a base in [2, n-2]
|
||||
let hash = U512::from_big_endian(sha3.finalize_reset().as_slice());
|
||||
let base = (hash % (*n - U512::from(4u32))) + U512::from(2u32);
|
||||
bases[base_count] = base;
|
||||
base_count += 1;
|
||||
|
||||
counter += U512::one();
|
||||
}
|
||||
|
||||
'witness: for base in bases {
|
||||
let mut x = mod_pow_impl(&base, &d, n);
|
||||
|
||||
if x == U512::one() || x == *n - U512::one() {
|
||||
continue 'witness;
|
||||
}
|
||||
|
||||
// Square r-1 times
|
||||
for _ in 0..r - 1 {
|
||||
x = mod_pow_impl(&x, &U512::from(2u32), n);
|
||||
if x == *n - U512::one() {
|
||||
continue 'witness;
|
||||
}
|
||||
if x == U512::one() {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
return false;
|
||||
}
|
||||
|
||||
true
|
||||
}
|
||||
|
||||
/// SHA3-512 over the big-endian bytes of input `U512`.
|
||||
fn sha3_512_impl(input: U512) -> U512 {
|
||||
use sha3::Digest;
|
||||
use sha3::Sha3_512;
|
||||
|
||||
let mut sha3 = Sha3_512::new();
|
||||
let bytes = input.to_big_endian();
|
||||
sha3.update(bytes);
|
||||
U512::from_big_endian(sha3.finalize().as_slice())
|
||||
}
|
||||
|
||||
#[cfg(test)]
|
||||
mod tests {
|
||||
use super::*;
|
||||
|
||||
// removed unused helper u512_from_hex
|
||||
|
||||
#[test]
|
||||
fn compat_distance_matches_context_distance() {
|
||||
// Synthetic header and nonce
|
||||
fn test_job_context_creation() {
|
||||
let header = [1u8; 32];
|
||||
let nonce = [2u8; 64];
|
||||
let difficulty = U512::from(1000u64);
|
||||
|
||||
let dist_compat = compat::get_nonce_distance(header, nonce);
|
||||
let ctx = JobContext::new(header, difficulty);
|
||||
|
||||
let ctx = JobContext::new(header, U512::from(123u64));
|
||||
let nonce_int = U512::from_big_endian(&nonce);
|
||||
let dist_ctx = distance_for_nonce(&ctx, nonce_int);
|
||||
|
||||
assert_eq!(dist_compat, dist_ctx);
|
||||
assert_eq!(ctx.header, header);
|
||||
assert_eq!(ctx.difficulty, difficulty);
|
||||
assert_eq!(ctx.target, U512::MAX / difficulty);
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn incremental_step_matches_pow_plus_one() {
|
||||
fn test_nonce_stepping() {
|
||||
let start = U512::from(100u64);
|
||||
let next = step_nonce(start);
|
||||
|
||||
assert_eq!(next, U512::from(101u64));
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn test_hash_computation() {
|
||||
let header = [1u8; 32];
|
||||
let difficulty = U512::from(1u64);
|
||||
let ctx = JobContext::new(header, difficulty);
|
||||
|
||||
let nonce = U512::from(123u64);
|
||||
let hash1 = hash_from_nonce(&ctx, nonce);
|
||||
let hash2 = hash_from_nonce(&ctx, nonce);
|
||||
|
||||
// Same input should produce same hash
|
||||
assert_eq!(hash1, hash2);
|
||||
|
||||
// Hash should not be zero for non-zero nonce
|
||||
assert_ne!(hash1, U512::zero());
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn test_different_nonces_different_hashes() {
|
||||
let header = [2u8; 32];
|
||||
let difficulty = U512::from(1u64);
|
||||
let ctx = JobContext::new(header, difficulty);
|
||||
|
||||
let nonce1 = U512::from(100u64);
|
||||
let nonce2 = U512::from(101u64);
|
||||
|
||||
let hash1 = hash_from_nonce(&ctx, nonce1);
|
||||
let hash2 = hash_from_nonce(&ctx, nonce2);
|
||||
|
||||
assert_ne!(hash1, hash2);
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn test_validity_check() {
|
||||
let header = [3u8; 32];
|
||||
let threshold = U512::from(99999u64);
|
||||
let ctx = JobContext::new(header, threshold);
|
||||
let easy_difficulty = U512::from(1u64);
|
||||
let ctx = JobContext::new(header, easy_difficulty);
|
||||
|
||||
let start = U512::from(1000u64);
|
||||
let y0 = init_worker_y0(&ctx, start);
|
||||
let nonce = U512::from(1u64);
|
||||
let (is_valid, hash) = is_valid_nonce(ctx.header, nonce.to_big_endian(), ctx.difficulty);
|
||||
|
||||
// Distance for start+1 computed two ways should match:
|
||||
// 1) Incremental: step once from y0
|
||||
// 2) Direct exponentiation with nonce = start+1
|
||||
let y1_inc = step_mul(&ctx, y0);
|
||||
let dist_inc = distance_from_y(&ctx, y1_inc);
|
||||
// With very easy difficulty, should be valid
|
||||
assert!(is_valid);
|
||||
assert_ne!(hash, U512::zero());
|
||||
|
||||
let direct = mod_pow_impl(
|
||||
&ctx.m,
|
||||
&ctx.header_int.saturating_add(start + U512::one()),
|
||||
&ctx.n,
|
||||
);
|
||||
let dist_direct = distance_from_y(&ctx, direct);
|
||||
|
||||
assert_eq!(dist_inc, dist_direct);
|
||||
// Verify hash is actually below target
|
||||
assert!(hash < ctx.target);
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn zero_nonce_is_zero_distance() {
|
||||
let header = [0xABu8; 32];
|
||||
let nonce = [0u8; 64];
|
||||
let d = compat::get_nonce_distance(header, nonce);
|
||||
assert_eq!(d, U512::zero());
|
||||
fn test_target_calculation() {
|
||||
let header = [4u8; 32];
|
||||
let difficulty = U512::from(256u64);
|
||||
let ctx = JobContext::new(header, difficulty);
|
||||
|
||||
let expected_target = U512::MAX / U512::from(256u64);
|
||||
assert_eq!(ctx.target, expected_target);
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn test_hash_matches_qpow_math() {
|
||||
// Test that our JobContext produces the same results as qpow_math directly
|
||||
let header = [1u8; 32];
|
||||
let nonce = U512::from(123u64);
|
||||
let difficulty = U512::from(1000u64);
|
||||
|
||||
let ctx = JobContext::new(header, difficulty);
|
||||
let hash_ctx = hash_from_nonce(&ctx, nonce);
|
||||
|
||||
let nonce_bytes = nonce.to_big_endian();
|
||||
let hash_direct = qpow_math::get_nonce_hash(header, nonce_bytes);
|
||||
|
||||
assert_eq!(hash_ctx, hash_direct);
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn test_mine_range_functionality() {
|
||||
let header = [5u8; 32];
|
||||
let difficulty = U512::from(1u64); // Very easy
|
||||
let ctx = JobContext::new(header, difficulty);
|
||||
|
||||
let start_nonce = U512::from(1u64);
|
||||
let result = mine_nonce_range(&ctx, start_nonce, 10);
|
||||
|
||||
// With very easy difficulty, should find a solution quickly
|
||||
if let Some((found_nonce, found_hash)) = result {
|
||||
assert!(found_nonce >= start_nonce);
|
||||
assert!(found_nonce < start_nonce + U512::from(10u64));
|
||||
assert!(found_hash < ctx.target);
|
||||
}
|
||||
// If no solution found, that's also valid behavior
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn test_hard_difficulty_no_solution() {
|
||||
let header = [6u8; 32];
|
||||
let very_hard_difficulty = U512::MAX; // Impossible difficulty
|
||||
let ctx = JobContext::new(header, very_hard_difficulty);
|
||||
|
||||
let start_nonce = U512::from(1u64);
|
||||
let result = mine_nonce_range(&ctx, start_nonce, 5);
|
||||
|
||||
// With impossible difficulty, should not find solution
|
||||
assert!(result.is_none());
|
||||
}
|
||||
}
|
||||
|
||||
@@ -28,7 +28,7 @@
|
||||
Description=Quantus External Miner Service
|
||||
Wants=network-online.target
|
||||
After=network-online.target
|
||||
Documentation=https://github.com/Quantus-Network/quantus-external-miner
|
||||
Documentation=https://github.com/Quantus-Network/quantus-miner
|
||||
|
||||
[Service]
|
||||
Type=simple
|
||||
|
||||
@@ -126,4 +126,4 @@ Operational tips
|
||||
- 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.
|
||||
|
||||
Support
|
||||
- Repository: https://github.com/Quantus-Network/quantus-external-miner
|
||||
- Repository: https://github.com/Quantus-Network/quantus-miner
|
||||
|
||||
48
src/lib.rs
48
src/lib.rs
@@ -1,48 +0,0 @@
|
||||
/*!
|
||||
Deprecated root library entrypoint.
|
||||
|
||||
This repository has been restructured into a Cargo workspace to provide clear
|
||||
separation of concerns and to enable aggressive optimization of the mining core.
|
||||
|
||||
Use the new workspace crates instead of the old root library:
|
||||
|
||||
- crates/miner-cli
|
||||
- The CLI binary for running the external miner service.
|
||||
- Examples:
|
||||
- cargo run -p miner-cli -- --port 9833
|
||||
- cargo run -p miner-cli -- --port 9833 --metrics-port 9900
|
||||
- cargo run -p miner-cli -- --workers 4
|
||||
|
||||
- crates/miner-service
|
||||
- The service layer: HTTP API (compatible with the node), job orchestration,
|
||||
and engine abstraction (CPU/GPU backends).
|
||||
|
||||
- crates/pow-core
|
||||
- The QPoW math core (local fork/superset of qpow-math) with a compatibility
|
||||
API and new optimized paths for future refactoring (e.g., precomputation,
|
||||
incremental evaluation, Montgomery multiplication).
|
||||
|
||||
- crates/engine-cpu
|
||||
- CPU mining engine(s) implementing the unified engine trait. Contains the
|
||||
baseline/reference engine and will host the optimized incremental/Montgomery
|
||||
engine.
|
||||
|
||||
- crates/engine-gpu-cuda (optional, scaffold)
|
||||
- Placeholder for a CUDA-based GPU engine.
|
||||
|
||||
- crates/engine-gpu-opencl (optional, scaffold)
|
||||
- Placeholder for an OpenCL-based GPU engine.
|
||||
|
||||
- crates/metrics (optional)
|
||||
- Prometheus metrics registry and optional HTTP exporter.
|
||||
- Metrics are toggled by the presence of the `--metrics-port` CLI parameter
|
||||
(when omitted, metrics are disabled entirely).
|
||||
|
||||
Notes:
|
||||
- The public HTTP API remains compatible with the node.
|
||||
- The old root library is intentionally left without exports to avoid conflicts.
|
||||
- Build the binary via the new CLI package:
|
||||
cargo build -p miner-cli --release
|
||||
- Run the service:
|
||||
cargo run -p miner-cli -- --port 9833 [--metrics-port 9900] [--workers N]
|
||||
*/
|
||||
33
src/main.rs
33
src/main.rs
@@ -1,33 +0,0 @@
|
||||
/*!
|
||||
Deprecated root binary entrypoint.
|
||||
|
||||
This repository has been restructured into a Cargo workspace.
|
||||
|
||||
Run the external miner using the new CLI binary:
|
||||
|
||||
- cargo run -p miner-cli -- [args...]
|
||||
- cargo build -p miner-cli --release
|
||||
|
||||
Examples:
|
||||
- cargo run -p miner-cli -- --port 9833
|
||||
- cargo run -p miner-cli -- --port 9833 --metrics-port 9900
|
||||
- cargo run -p miner-cli -- --workers 4
|
||||
|
||||
See crates/miner-cli for the active entrypoint and crates/miner-service for the service layer.
|
||||
*/
|
||||
|
||||
fn main() {
|
||||
eprintln!(
|
||||
"\n[DEPRECATED] This entrypoint has moved.\n\
|
||||
Use the new CLI binary in the workspace instead:\n\
|
||||
\n\
|
||||
- cargo run -p miner-cli -- [args...]\n\
|
||||
- cargo build -p miner-cli --release\n\
|
||||
\n\
|
||||
Examples:\n\
|
||||
- cargo run -p miner-cli -- --port 9833\n\
|
||||
- cargo run -p miner-cli -- --port 9833 --metrics-port 9900\n\
|
||||
- cargo run -p miner-cli -- --workers 4\n"
|
||||
);
|
||||
std::process::exit(1);
|
||||
}
|
||||
Reference in New Issue
Block a user