Compare commits
| Author | SHA1 | Date | |
|---|---|---|---|
 Quantus Network CIandczareko
|
a6660c5e13 | ||
|
illuzen
|
85006dd3a7 | ||
|
illuzen
|
0a96d2f6b7 | ||
|
illuzen
|
aa1204bab9 |
Generated
+190
-43
@@ -460,6 +460,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,7 +477,7 @@ dependencies = [
|
||||
|
||||
[[package]]
|
||||
name = "engine-cpu"
|
||||
version = "0.2.1"
|
||||
version = "0.3.0"
|
||||
dependencies = [
|
||||
"anyhow",
|
||||
"hex",
|
||||
@@ -483,7 +489,7 @@ dependencies = [
|
||||
|
||||
[[package]]
|
||||
name = "engine-gpu-cuda"
|
||||
version = "0.2.1"
|
||||
version = "0.3.0"
|
||||
dependencies = [
|
||||
"anyhow",
|
||||
"cust",
|
||||
@@ -495,14 +501,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 = "0.3.0"
|
||||
dependencies = [
|
||||
"anyhow",
|
||||
"log",
|
||||
@@ -514,7 +520,7 @@ dependencies = [
|
||||
|
||||
[[package]]
|
||||
name = "engine-montgomery"
|
||||
version = "0.2.1"
|
||||
version = "0.3.0"
|
||||
dependencies = [
|
||||
"crypto-bigint",
|
||||
"engine-cpu",
|
||||
@@ -523,7 +529,6 @@ dependencies = [
|
||||
"num-traits 0.2.19",
|
||||
"pow-core",
|
||||
"primitive-types",
|
||||
"sha3",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
@@ -1058,6 +1063,15 @@ version = "1.70.1"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "7943c866cc5cd64cbc25b2e01621d07fa8eb2a1a23160ee81ce38704e97b8ecf"
|
||||
|
||||
[[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 +1088,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 +1136,7 @@ checksum = "32a282da65faaf38286cf3be983213fcf1d2e2a58700e808f83f4ea9a4804bc0"
|
||||
|
||||
[[package]]
|
||||
name = "metrics"
|
||||
version = "0.2.1"
|
||||
version = "0.3.0"
|
||||
dependencies = [
|
||||
"anyhow",
|
||||
"log",
|
||||
@@ -1161,7 +1166,7 @@ dependencies = [
|
||||
|
||||
[[package]]
|
||||
name = "miner-cli"
|
||||
version = "0.2.1"
|
||||
version = "0.3.0"
|
||||
dependencies = [
|
||||
"clap",
|
||||
"env_logger",
|
||||
@@ -1172,7 +1177,7 @@ dependencies = [
|
||||
|
||||
[[package]]
|
||||
name = "miner-service"
|
||||
version = "0.2.1"
|
||||
version = "0.3.0"
|
||||
dependencies = [
|
||||
"anyhow",
|
||||
"crossbeam-channel",
|
||||
@@ -1197,7 +1202,7 @@ dependencies = [
|
||||
|
||||
[[package]]
|
||||
name = "miner-telemetry"
|
||||
version = "0.2.1"
|
||||
version = "0.3.0"
|
||||
dependencies = [
|
||||
"anyhow",
|
||||
"futures 0.3.31",
|
||||
@@ -1378,6 +1383,122 @@ version = "1.70.1"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "a4895175b425cb1f87721b59f0f286c2092bd4af812243672510e1ac53e2e0ad"
|
||||
|
||||
[[package]]
|
||||
name = "p3-dft"
|
||||
version = "0.3.0"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "e3b2764a3982d22d62aa933c8de6f9d71d8a474c9110b69e675dea1887bdeffc"
|
||||
dependencies = [
|
||||
"itertools",
|
||||
"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",
|
||||
"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",
|
||||
"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",
|
||||
"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 +1550,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"
|
||||
@@ -1478,7 +1605,7 @@ dependencies = [
|
||||
|
||||
[[package]]
|
||||
name = "pow-core"
|
||||
version = "0.2.1"
|
||||
version = "0.3.0"
|
||||
dependencies = [
|
||||
"anyhow",
|
||||
"hex",
|
||||
@@ -1486,8 +1613,7 @@ dependencies = [
|
||||
"num-bigint",
|
||||
"num-traits 0.2.19",
|
||||
"primitive-types",
|
||||
"sha2",
|
||||
"sha3",
|
||||
"qp-poseidon-core",
|
||||
"thiserror",
|
||||
]
|
||||
|
||||
@@ -1543,6 +1669,20 @@ dependencies = [
|
||||
"thiserror",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "qp-poseidon-core"
|
||||
version = "0.9.5"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "ec326fc2631a929de09a38af2613a3db5230882c12a2f68205693ec632751e8b"
|
||||
dependencies = [
|
||||
"p3-field",
|
||||
"p3-goldilocks",
|
||||
"p3-poseidon2",
|
||||
"p3-symmetric",
|
||||
"rand 0.9.2",
|
||||
"rand_chacha 0.9.0",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "quantus-miner-api"
|
||||
version = "0.0.3"
|
||||
@@ -1873,27 +2013,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 +2078,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"
|
||||
@@ -2175,9 +2300,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 +2324,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"
|
||||
|
||||
+2
-3
@@ -17,7 +17,7 @@ resolver = "2"
|
||||
edition = "2021"
|
||||
authors = ["Quantus Network"]
|
||||
description = "Quantus External Miner Workspace"
|
||||
version = "0.2.1"
|
||||
version = "0.3.0"
|
||||
|
||||
[workspace.dependencies]
|
||||
anyhow = "1"
|
||||
@@ -31,12 +31,11 @@ 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 = "0.9.5", 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"
|
||||
|
||||
@@ -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" }
|
||||
|
||||
@@ -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.
|
||||
@@ -19,7 +19,7 @@
|
||||
//! This crate now includes a host-side launcher for a minimal CUDA kernel
|
||||
//! (see `src/kernels/qpow_kernel.cu`) that performs 512-bit Montgomery
|
||||
//! multiplication on-device and returns normalized y values. For G1 bring-up,
|
||||
//! we perform SHA3 and threshold checks on the host for parity against CPU.
|
||||
//! we perform Poseidon2 and threshold checks on the host for parity against CPU.
|
||||
//!
|
||||
//! Behavior:
|
||||
//! - If compiled without the `cuda` feature or if CUDA init/launch fails,
|
||||
@@ -462,7 +462,7 @@ impl CudaEngine {
|
||||
total_bytes
|
||||
);
|
||||
|
||||
// Branch kernel launch by mode: G2 (device SHA3 + early-exit) vs G1 (return y values)
|
||||
// Branch kernel launch by mode: G2 (device Poseidon2 + early-exit) vs G1 (return y values)
|
||||
if is_g2 {
|
||||
// Compute full-precision bounds for accounting (use precomputed)
|
||||
let rem_be = rem_be_all;
|
||||
@@ -1050,7 +1050,7 @@ impl CudaEngine {
|
||||
|
||||
// G2 handled in the launch branch above; proceed with G1 copy-back path below.
|
||||
if !is_g2 {
|
||||
// G1 path (host SHA3) — Copy back results (optional pinned host buffer + async copy)
|
||||
// G1 path (host Poseidon2) — Copy back results (optional pinned host buffer + async copy)
|
||||
let use_pinned = std::env::var("MINER_CUDA_PINNED")
|
||||
.ok()
|
||||
.map(|v| v == "1" || v.eq_ignore_ascii_case("true"))
|
||||
@@ -1086,7 +1086,7 @@ impl CudaEngine {
|
||||
std::sync::Arc::new(y_out_host)
|
||||
};
|
||||
|
||||
// Parallel SHA3 over GPU results; compute earliest valid index (if any)
|
||||
// Parallel Poseidon2 over GPU results; compute earliest valid index (if any)
|
||||
let t_sha_start = std::time::Instant::now();
|
||||
|
||||
// Bound hashing work to remaining range in this chunk (full-precision)
|
||||
@@ -1126,7 +1126,7 @@ impl CudaEngine {
|
||||
.map(|n| n.get())
|
||||
.unwrap_or(1)
|
||||
});
|
||||
log::info!(target: "miner", "Host SHA3: threads={hash_threads}, total_elems={total_to_hash}");
|
||||
log::info!(target: "miner", "Host Poseidon2: threads={hash_threads}, total_elems={total_to_hash}");
|
||||
|
||||
// Share outputs across worker threads (Vec or LockedBuffer depending on MINER_CUDA_PINNED)
|
||||
let y_shared = y_host_arc.clone();
|
||||
@@ -1178,9 +1178,9 @@ impl CudaEngine {
|
||||
}
|
||||
});
|
||||
|
||||
// SHA3 (host) timing
|
||||
let sha3_ms = t_sha_start.elapsed().as_millis();
|
||||
log::info!(target: "miner", "SHA3 (host) OK: sha3_ms={sha3_ms}");
|
||||
// Poseidon2 (host) timing
|
||||
let poseidon2_ms = t_sha_start.elapsed().as_millis();
|
||||
log::info!(target: "miner", "Poseidon2 (host) OK: poseidon2_ms={poseidon2_ms}");
|
||||
|
||||
// Outcomes and accounting
|
||||
let total_elems_u64 = total_to_hash as u64;
|
||||
|
||||
@@ -21,7 +21,7 @@ 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 }
|
||||
|
||||
|
||||
@@ -11,7 +11,7 @@ The engine mirrors `cpu-fast` control flow and metrics emission so that apples-t
|
||||
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.
|
||||
- Direct Poseidon2 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.
|
||||
@@ -26,7 +26,7 @@ QPoW distance per nonce is computed as:
|
||||
- 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.
|
||||
2) Distance = target XOR Poseidon2_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).
|
||||
|
||||
@@ -64,8 +64,8 @@ This yields a single multiply+reduce per nonce with a branchless inner loop (exc
|
||||
## 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.
|
||||
- Before Poseidon2, we convert the residue `ŷ` back to the normal domain via `mont_mul(ŷ, 1)` and serialize as big-endian 64 bytes.
|
||||
- We reuse a single `Poseidon2Core` hasher per search call and `hash_512()` each iteration to reduce construction overhead.
|
||||
|
||||
This preserves consensus behavior while avoiding unnecessary big-int allocations or conversions.
|
||||
|
||||
@@ -148,10 +148,10 @@ We recommend running the property tests on machines with and without BMI2/ADX an
|
||||
|
||||
## Performance Notes
|
||||
|
||||
- Relative gains depend on how much time the miner spends in SHA3 vs modular multiply.
|
||||
- Relative gains depend on how much time the miner spends in Poseidon2 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.
|
||||
- BMI2+ADX typically produces the highest gains on the multiply itself (often 1.5–2.0×), with end-to-end uplift bounded by Poseidon2 share per nonce.
|
||||
- aarch64 UMULH/ADCS brings similar relative gains on Apple Silicon and other ARM64 platforms.
|
||||
|
||||
To minimize orchestration overhead in the service:
|
||||
@@ -181,7 +181,7 @@ To minimize orchestration overhead in the service:
|
||||
|
||||
- [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] Direct Poseidon2 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
|
||||
|
||||
@@ -19,7 +19,7 @@
|
||||
//! - 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.
|
||||
//! - Convert out of Montgomery before Poseidon2-512 distance computation.
|
||||
//!
|
||||
//! Important: We intentionally keep emissions (hash_count increments, control flow)
|
||||
//! identical to `cpu-fast` so metrics are directly comparable when pitting
|
||||
|
||||
@@ -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,7 @@ 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 }
|
||||
thiserror = { workspace = true }
|
||||
anyhow = { workspace = true }
|
||||
|
||||
|
||||
+32
-39
@@ -11,12 +11,13 @@
|
||||
//
|
||||
// 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.
|
||||
// - Future work will gate accelerated paths (e.g., Montgomery, SIMD Poseidon2) behind features.
|
||||
|
||||
extern crate alloc;
|
||||
|
||||
use core::ops::BitXor;
|
||||
use primitive_types::U512;
|
||||
use qp_poseidon_core::Poseidon2Core;
|
||||
|
||||
#[cfg(feature = "std")]
|
||||
use log::{debug, error};
|
||||
@@ -59,7 +60,7 @@ pub mod compat {
|
||||
/// 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.
|
||||
/// where (m, n) are derived deterministically from the header, and H is Poseidon2-512.
|
||||
pub fn get_nonce_distance(header: [u8; 32], nonce: [u8; 64]) -> U512 {
|
||||
super::get_nonce_distance_impl(header, nonce)
|
||||
}
|
||||
@@ -79,7 +80,7 @@ pub mod compat {
|
||||
super::is_prime_impl(n)
|
||||
}
|
||||
|
||||
/// Apply the reference "hash-to-group" function then SHA3-512.
|
||||
/// Apply the reference "hash-to-group" function then Poseidon2-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)
|
||||
}
|
||||
@@ -94,9 +95,9 @@ pub mod compat {
|
||||
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)
|
||||
/// Poseidon2-512 over the big-endian bytes of input U512.
|
||||
pub fn poseidon2_512(input: U512) -> U512 {
|
||||
super::poseidon2_512_impl(input)
|
||||
}
|
||||
}
|
||||
|
||||
@@ -150,9 +151,9 @@ pub fn step_mul(ctx: &JobContext, y: U512) -> U512 {
|
||||
}
|
||||
|
||||
/// Compute distance for the current y:
|
||||
/// distance = target XOR SHA3_512(y)
|
||||
/// distance = target XOR Poseidon2_512(y)
|
||||
pub fn distance_from_y(ctx: &JobContext, y: U512) -> U512 {
|
||||
let hashed = sha3_512_impl(y);
|
||||
let hashed = poseidon2_512_impl(y);
|
||||
ctx.target.bitxor(hashed)
|
||||
}
|
||||
|
||||
@@ -190,26 +191,24 @@ fn get_nonce_distance_impl(header: [u8; 32], nonce: [u8; 64]) -> U512 {
|
||||
|
||||
/// 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:
|
||||
/// - m: 256-bit derived via Poseidon2-256(header)
|
||||
/// - n: 512-bit derived via Poseidon2-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;
|
||||
let poseidon = Poseidon2Core::new();
|
||||
|
||||
// m from SHA2-256
|
||||
let mut sha256 = Sha256::new();
|
||||
sha256.update(header);
|
||||
let m = U512::from_big_endian(sha256.finalize().as_slice());
|
||||
// m from Poseidon2-256
|
||||
let m_bytes = poseidon.hash_no_pad_bytes(header);
|
||||
let m = U512::from_big_endian(&m_bytes);
|
||||
|
||||
// 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());
|
||||
// initial n from Poseidon2-512
|
||||
let mut n_bytes = poseidon.hash_512(&m_bytes);
|
||||
let mut n = U512::from_big_endian(&n_bytes);
|
||||
|
||||
// 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);
|
||||
n_bytes = poseidon.hash_512(&n_bytes);
|
||||
n = U512::from_big_endian(&n_bytes);
|
||||
}
|
||||
|
||||
(m, n)
|
||||
@@ -229,12 +228,12 @@ fn is_coprime_impl(a: &U512, b: &U512) -> bool {
|
||||
x == U512::one()
|
||||
}
|
||||
|
||||
/// Hash-to-group then SHA3-512.
|
||||
/// Hash-to-group then Poseidon2-512.
|
||||
///
|
||||
/// Note: The reference calls `hash_to_group_bigint` followed by an additional SHA3-512.
|
||||
/// Note: The reference calls `hash_to_group_bigint` followed by an additional Poseidon2-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)
|
||||
poseidon2_512_impl(result)
|
||||
}
|
||||
|
||||
/// Reference hash-to-group big-integer function (no chunk splitting).
|
||||
@@ -291,11 +290,9 @@ fn mod_mul_impl(a: &U512, b: &U512, modulus: &U512) -> U512 {
|
||||
|
||||
/// Miller–Rabin primality test.
|
||||
///
|
||||
/// Deterministically selects k=32 bases hashed from `n` using SHA3-512 to
|
||||
/// Deterministically selects k=32 bases hashed from `n` using Poseidon2-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;
|
||||
}
|
||||
@@ -314,10 +311,10 @@ fn is_prime_impl(n: &U512) -> bool {
|
||||
r += 1;
|
||||
}
|
||||
|
||||
// Generate test bases deterministically from n using SHA3
|
||||
// Generate test bases deterministically from n using Poseidon2
|
||||
let mut bases = [U512::zero(); 32];
|
||||
let mut base_count = 0;
|
||||
let mut sha3 = Sha3_512::new();
|
||||
let poseidon = Poseidon2Core::new();
|
||||
let mut counter = U512::zero();
|
||||
|
||||
while base_count < 32 {
|
||||
@@ -329,10 +326,10 @@ fn is_prime_impl(n: &U512) -> bool {
|
||||
bytes[..64].copy_from_slice(&n_bytes);
|
||||
bytes[64..128].copy_from_slice(&counter_bytes);
|
||||
|
||||
sha3.update(bytes);
|
||||
let poseidon_bytes = poseidon.hash_512(&bytes);
|
||||
|
||||
// Use the hash to generate a base in [2, n-2]
|
||||
let hash = U512::from_big_endian(sha3.finalize_reset().as_slice());
|
||||
let hash = U512::from_big_endian(&poseidon_bytes);
|
||||
let base = (hash % (*n - U512::from(4u32))) + U512::from(2u32);
|
||||
bases[base_count] = base;
|
||||
base_count += 1;
|
||||
@@ -363,15 +360,11 @@ fn is_prime_impl(n: &U512) -> bool {
|
||||
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();
|
||||
/// Poseidon2-512 over the big-endian bytes of input `U512`.
|
||||
fn poseidon2_512_impl(input: U512) -> U512 {
|
||||
let poseidon = Poseidon2Core::new();
|
||||
let bytes = input.to_big_endian();
|
||||
sha3.update(bytes);
|
||||
U512::from_big_endian(sha3.finalize().as_slice())
|
||||
U512::from_big_endian(&poseidon.hash_512(&bytes))
|
||||
}
|
||||
|
||||
#[cfg(test)]
|
||||
|
||||
Reference in New Issue
Block a user