engine-cuda: native CUDA mining engine behind MinerEngine
All checks were successful
ci / fmt (pull_request) Successful in 20s
bench / build (pull_request) Successful in 1m15s
ci / clippy (pull_request) Successful in 1m40s
ci / doc (pull_request) Successful in 2m6s
bench / measure (pull_request) Successful in 2m51s
ci / test (pull_request) Successful in 7m45s
All checks were successful
ci / fmt (pull_request) Successful in 20s
bench / build (pull_request) Successful in 1m15s
ci / clippy (pull_request) Successful in 1m40s
ci / doc (pull_request) Successful in 2m6s
bench / measure (pull_request) Successful in 2m51s
ci / test (pull_request) Successful in 7m45s
quantus/miner#3. New crate, nothing in engine-gpu touched. - kernels/mining.cu: the same host contract and sponge schedule as mining_u64.wgsl (host midstate, low-256-bit nonce increments, lazy second squeeze), bit-exact with pow_core. The field multiply is __umul64hi plus a plain multiply and one reduction, instead of four 32-bit partials with carry reconstruction. - build.rs: generates poseidon2_constants.cuh from qp-poseidon-constants (a hash change at origin is a dependency bump, never a hand-copied table) and, when nvcc is present, compiles one fat binary with cubins for sm_86/89/120 plus compute_120 PTX. Without nvcc it writes an empty image, the crate still builds, and CudaEngine::try_new fails with a clear message so the miner falls back to wgpu. - lib.rs: CudaEngine with the same batch loop, cancellation, thread-local device assignment, metrics (kernel label "cuda") and stale accounting as engine-gpu. cudarc 0.19 with dynamic loading; no link-time CUDA dependency. - miner-service: resolve_gpu_configuration tries CUDA first unless --gpu-engine wgpu; --gpu-engine cuda makes its absence an error. miner-cli gains the flag on serve and benchmark (MINER_GPU_ENGINE). - bench-harness: --engine auto|cuda|wgpu. - deploy.yaml: matrix rows carry the kernel validate expects; a wgpu fallback on a CUDA host now fails the deploy instead of passing at a fraction of the hashrate. Measured on quadbrat (RTX 3060, 130 W, miner paused, 2 x 10 s windows, parity 5/5 verified against CPU): wgpu 37.3 MH/s, CUDA 59.5 MH/s (1.59x). Grid sizing 8192 threads/SM or more is flat; 2048 loses 10%. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01CBgs2nSi4H2mdh8kD8vMX5
This commit is contained in:
@@ -82,15 +82,21 @@ jobs:
|
||||
# Policy until the mainnet call (quantus/miner#1): benjy and quadbrat
|
||||
# mine testnet continuously; beast (2x RTX 5090) serves inference and
|
||||
# does not mine. Enabling it is uncommenting the entry.
|
||||
# `kernel` is the kernel id validate expects on the per-device metric:
|
||||
# cuda (engine-cuda, #3) on every NVIDIA host once the build carries
|
||||
# the fat binary; u64 would mean the miner silently fell back to wgpu.
|
||||
- host: benjy.hanzalova.internal
|
||||
node: bob.hanzalova.internal
|
||||
gpu_devices: "1" # 1x RTX 4090 (sm_89); reference card for #2
|
||||
kernel: cuda
|
||||
- host: quadbrat.hanzalova.internal
|
||||
node: bob.hanzalova.internal
|
||||
gpu_devices: "1" # 1x RTX 3060 (sm_86)
|
||||
kernel: cuda
|
||||
# - host: beast.hanzalova.internal
|
||||
# node: bob.hanzalova.internal
|
||||
# gpu_devices: "2" # 2x RTX 5090 (sm_120)
|
||||
# kernel: cuda
|
||||
steps:
|
||||
- uses: actions/checkout@v4
|
||||
- uses: actions/download-artifact@v3
|
||||
@@ -314,6 +320,15 @@ jobs:
|
||||
elif [ "${{ github.event.inputs.mode }}" != validate ]; then
|
||||
echo " miner_build_info does not carry commit ${{ github.sha }}" >&2; fail=1
|
||||
fi
|
||||
# The kernel actually running. A wgpu fallback on a host that should
|
||||
# run CUDA passes every other check at a fraction of the hashrate.
|
||||
if grep -q "^miner_device_hashes_total{.*kernel=\"${{ matrix.kernel }}\"" <<<"$m"; then
|
||||
echo " ok kernel ${{ matrix.kernel }} on device 0"
|
||||
else
|
||||
echo " expected kernel ${{ matrix.kernel }}, exported series:" >&2
|
||||
grep "^miner_device_hashes_total" <<<"$m" >&2 || echo " (none yet)" >&2
|
||||
if [ "${{ github.event.inputs.mode }}" != validate ]; then fail=1; fi
|
||||
fi
|
||||
|
||||
deadline=$((SECONDS + 120))
|
||||
h1=""; h2=""; ok=0
|
||||
|
||||
46
Cargo.lock
generated
46
Cargo.lock
generated
@@ -107,7 +107,7 @@ version = "0.38.0+1.3.281"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "0bb44936d800fea8f016d7f2311c6a4f97aebd5dc86f09906139ec848cf3a46f"
|
||||
dependencies = [
|
||||
"libloading",
|
||||
"libloading 0.8.9",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
@@ -128,6 +128,7 @@ version = "4.0.2"
|
||||
dependencies = [
|
||||
"clap",
|
||||
"engine-cpu",
|
||||
"engine-cuda",
|
||||
"engine-gpu",
|
||||
"env_logger",
|
||||
"log",
|
||||
@@ -540,6 +541,15 @@ dependencies = [
|
||||
"typenum",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "cudarc"
|
||||
version = "0.19.9"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "804764d10e844da09765a7b2ca9641a0851523d1702efb0d7299d73e31b86e80"
|
||||
dependencies = [
|
||||
"libloading 0.9.0",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "data-encoding"
|
||||
version = "2.10.0"
|
||||
@@ -602,6 +612,19 @@ dependencies = [
|
||||
"rand 0.9.2",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "engine-cuda"
|
||||
version = "4.0.2"
|
||||
dependencies = [
|
||||
"cudarc",
|
||||
"engine-cpu",
|
||||
"log",
|
||||
"metrics",
|
||||
"pow-core",
|
||||
"primitive-types 0.13.1",
|
||||
"qp-poseidon-constants",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "engine-gpu"
|
||||
version = "4.0.2"
|
||||
@@ -653,7 +676,7 @@ source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "39cab71617ae0d63f51a36d69f866391735b51691dbda63cf6f96d042b63efeb"
|
||||
dependencies = [
|
||||
"libc",
|
||||
"windows-sys 0.52.0",
|
||||
"windows-sys 0.61.2",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
@@ -1293,7 +1316,7 @@ checksum = "3640c1c38b8e4e43584d8df18be5fc6b0aa314ce6ebf51b53313d4306cca8e46"
|
||||
dependencies = [
|
||||
"hermit-abi",
|
||||
"libc",
|
||||
"windows-sys 0.52.0",
|
||||
"windows-sys 0.61.2",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
@@ -1390,7 +1413,7 @@ source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "6aae1df220ece3c0ada96b8153459b67eebe9ae9212258bb0134ae60416fdf76"
|
||||
dependencies = [
|
||||
"libc",
|
||||
"libloading",
|
||||
"libloading 0.8.9",
|
||||
"pkg-config",
|
||||
]
|
||||
|
||||
@@ -1428,6 +1451,16 @@ dependencies = [
|
||||
"windows-link",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "libloading"
|
||||
version = "0.9.0"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "754ca22de805bb5744484a5b151a9e1a8e837d5dc232c2d7d8c2e3492edc8b60"
|
||||
dependencies = [
|
||||
"cfg-if",
|
||||
"windows-link",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
name = "libm"
|
||||
version = "0.2.16"
|
||||
@@ -1544,6 +1577,7 @@ dependencies = [
|
||||
"anyhow",
|
||||
"crossbeam-channel",
|
||||
"engine-cpu",
|
||||
"engine-cuda",
|
||||
"engine-gpu",
|
||||
"getrandom 0.2.17",
|
||||
"hex",
|
||||
@@ -3719,7 +3753,7 @@ dependencies = [
|
||||
"js-sys",
|
||||
"khronos-egl",
|
||||
"libc",
|
||||
"libloading",
|
||||
"libloading 0.8.9",
|
||||
"log",
|
||||
"metal",
|
||||
"naga",
|
||||
@@ -3779,7 +3813,7 @@ version = "0.1.11"
|
||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||
checksum = "c2a7b1c03c876122aa43f3020e6c3c3ee5c05081c9a00739faf7503aeba10d22"
|
||||
dependencies = [
|
||||
"windows-sys 0.48.0",
|
||||
"windows-sys 0.61.2",
|
||||
]
|
||||
|
||||
[[package]]
|
||||
|
||||
@@ -2,6 +2,7 @@
|
||||
members = [
|
||||
"crates/bench-harness", # lair: quantus/miner#2
|
||||
"crates/engine-cpu",
|
||||
"crates/engine-cuda", # lair: quantus/miner#3
|
||||
"crates/engine-gpu",
|
||||
"crates/metrics",
|
||||
"crates/miner-cli",
|
||||
|
||||
@@ -12,6 +12,7 @@ path = "src/main.rs"
|
||||
[dependencies]
|
||||
engine-cpu = { path = "../engine-cpu" }
|
||||
engine-gpu = { path = "../engine-gpu" }
|
||||
engine-cuda = { path = "../engine-cuda" }
|
||||
pow-core = { path = "../pow-core" }
|
||||
primitive-types = { workspace = true }
|
||||
clap = { workspace = true, features = ["derive", "env"] }
|
||||
|
||||
@@ -11,6 +11,7 @@
|
||||
|
||||
use clap::Parser;
|
||||
use engine_cpu::{AtomicBoolCancelCheck, EngineStatus, MinerEngine, Range};
|
||||
use engine_cuda::CudaEngine;
|
||||
use engine_gpu::GpuEngine;
|
||||
use primitive_types::U512;
|
||||
use rand::RngCore;
|
||||
@@ -40,6 +41,11 @@ struct Args {
|
||||
#[arg(long, default_value_t = 1_000_000)]
|
||||
batch_size: u32,
|
||||
|
||||
/// GPU engine: auto (cuda if this binary carries a kernel and a driver is
|
||||
/// present, else wgpu), cuda, or wgpu.
|
||||
#[arg(long, default_value = "auto")]
|
||||
engine: String,
|
||||
|
||||
/// Worker threads. The engine assigns threads to devices round-robin, so
|
||||
/// on a multi-card host N threads measure N cards. More threads than cards
|
||||
/// is allowed on purpose: two threads on one card overlap one thread's
|
||||
@@ -128,6 +134,47 @@ struct Record {
|
||||
parity: Option<Parity>,
|
||||
}
|
||||
|
||||
/// Either GPU engine behind the same trait; the harness drives them identically.
|
||||
enum Engine {
|
||||
Wgpu(GpuEngine),
|
||||
Cuda(CudaEngine),
|
||||
}
|
||||
|
||||
impl Engine {
|
||||
fn open(kind: &str, batch_size: u32) -> Engine {
|
||||
match kind {
|
||||
"cuda" => {
|
||||
Engine::Cuda(CudaEngine::try_new(batch_size, 0).expect("CUDA engine init failed"))
|
||||
}
|
||||
"wgpu" => Engine::Wgpu(
|
||||
GpuEngine::try_new(batch_size, 0, false).expect("GPU engine init failed"),
|
||||
),
|
||||
"auto" => match CudaEngine::try_new(batch_size, 0) {
|
||||
Ok(e) => Engine::Cuda(e),
|
||||
Err(e) => {
|
||||
log::info!("CUDA engine unavailable ({e}); using wgpu");
|
||||
Engine::Wgpu(
|
||||
GpuEngine::try_new(batch_size, 0, false).expect("GPU engine init failed"),
|
||||
)
|
||||
}
|
||||
},
|
||||
other => panic!("unknown --engine {other}; use auto, cuda or wgpu"),
|
||||
}
|
||||
}
|
||||
fn device_count(&self) -> usize {
|
||||
match self {
|
||||
Engine::Wgpu(e) => e.device_count(),
|
||||
Engine::Cuda(e) => e.device_count(),
|
||||
}
|
||||
}
|
||||
fn as_dyn(&self) -> &dyn MinerEngine {
|
||||
match self {
|
||||
Engine::Wgpu(e) => e,
|
||||
Engine::Cuda(e) => e,
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
fn nvidia_smi() -> Vec<GpuState> {
|
||||
let out = Command::new("nvidia-smi")
|
||||
.args([
|
||||
@@ -192,7 +239,7 @@ fn median(xs: &[f64]) -> f64 {
|
||||
/// returning hashes and elapsed seconds. The difficulty is high enough that
|
||||
/// a solution is effectively impossible, matching production where the second
|
||||
/// squeeze is almost never taken.
|
||||
fn window(engine: &GpuEngine, secs: u64, rng: &mut impl RngCore) -> (u64, f64) {
|
||||
fn window(engine: &dyn MinerEngine, secs: u64, rng: &mut impl RngCore) -> (u64, f64) {
|
||||
let mut header = [0u8; 32];
|
||||
rng.fill_bytes(&mut header);
|
||||
let ctx = engine.prepare_context(header, U512::from(1u64) << 200);
|
||||
@@ -231,7 +278,7 @@ fn window(engine: &GpuEngine, secs: u64, rng: &mut impl RngCore) -> (u64, f64) {
|
||||
(hashes, elapsed)
|
||||
}
|
||||
|
||||
fn parity(engine: &GpuEngine, jobs: usize) -> Parity {
|
||||
fn parity(engine: &dyn MinerEngine, jobs: usize) -> Parity {
|
||||
// Same shape as engine-gpu's gpu_cpu_parity example, kept in lockstep with it.
|
||||
let mut rng = rand::rng();
|
||||
let cancel_flag = AtomicBool::new(false);
|
||||
@@ -305,8 +352,7 @@ fn main() {
|
||||
}
|
||||
}
|
||||
|
||||
let engine =
|
||||
Arc::new(GpuEngine::try_new(args.batch_size, 0, false).expect("GPU engine init failed"));
|
||||
let engine = Arc::new(Engine::open(&args.engine, args.batch_size));
|
||||
let devices = engine.device_count();
|
||||
if args.workers > devices {
|
||||
log::warn!(
|
||||
@@ -317,7 +363,7 @@ fn main() {
|
||||
}
|
||||
log::info!(
|
||||
"engine {} with {devices} device(s); measuring {} worker(s), {} x {}s windows after {}s warm-up, batch {}",
|
||||
engine.name(),
|
||||
engine.as_dyn().name(),
|
||||
args.workers,
|
||||
args.runs,
|
||||
args.duration_secs,
|
||||
@@ -332,11 +378,11 @@ fn main() {
|
||||
handles.push(std::thread::spawn(move || {
|
||||
let mut rng = rand::rng();
|
||||
if warm > 0 {
|
||||
let _ = window(&engine, warm, &mut rng);
|
||||
let _ = window(engine.as_dyn(), warm, &mut rng);
|
||||
}
|
||||
let mut windows = Vec::with_capacity(runs);
|
||||
for i in 0..runs {
|
||||
let (h, s) = window(&engine, dur, &mut rng);
|
||||
let (h, s) = window(engine.as_dyn(), dur, &mut rng);
|
||||
let mhs = h as f64 / s / 1e6;
|
||||
log::info!("worker {w} window {i}: {h} hashes in {s:.2}s = {mhs:.2} MH/s");
|
||||
windows.push(mhs);
|
||||
@@ -362,7 +408,7 @@ fn main() {
|
||||
let gpu_after = nvidia_smi();
|
||||
|
||||
let parity_result = if args.parity_jobs > 0 {
|
||||
Some(parity(&engine, args.parity_jobs))
|
||||
Some(parity(engine.as_dyn(), args.parity_jobs))
|
||||
} else {
|
||||
None
|
||||
};
|
||||
@@ -376,7 +422,7 @@ fn main() {
|
||||
host: hostname(),
|
||||
commit: args.commit.clone(),
|
||||
label: args.label.clone(),
|
||||
engine: engine.name().to_string(),
|
||||
engine: engine.as_dyn().name().to_string(),
|
||||
batch_size: args.batch_size,
|
||||
duration_secs: args.duration_secs,
|
||||
runs: args.runs,
|
||||
|
||||
17
crates/engine-cuda/Cargo.toml
Normal file
17
crates/engine-cuda/Cargo.toml
Normal file
@@ -0,0 +1,17 @@
|
||||
[package]
|
||||
name = "engine-cuda"
|
||||
version.workspace = true
|
||||
edition.workspace = true
|
||||
publish = false
|
||||
description = "lair: native CUDA mining engine behind MinerEngine (quantus/miner#3)"
|
||||
|
||||
[dependencies]
|
||||
engine-cpu = { path = "../engine-cpu" }
|
||||
pow-core = { path = "../pow-core" }
|
||||
metrics = { path = "../metrics" }
|
||||
primitive-types = { workspace = true }
|
||||
log = { workspace = true }
|
||||
cudarc = { version = "0.19", default-features = false, features = ["driver", "nvrtc", "dynamic-loading", "cuda-13000"] }
|
||||
|
||||
[build-dependencies]
|
||||
qp-poseidon-constants = "1.1"
|
||||
189
crates/engine-cuda/build.rs
Normal file
189
crates/engine-cuda/build.rs
Normal file
@@ -0,0 +1,189 @@
|
||||
//! lair: build the CUDA mining kernel into a fat binary (quantus/miner#3, #8).
|
||||
//!
|
||||
//! 1. Generate `poseidon2_constants.cuh` from `qp-poseidon-constants`, so a
|
||||
//! change to the hash at origin reaches the kernel as a dependency bump and
|
||||
//! never as a hand-copied table (rule from quantus/miner#1).
|
||||
//! 2. If `nvcc` is available, compile `src/kernels/mining.cu` once with one
|
||||
//! cubin per fleet architecture plus PTX for the newest, into
|
||||
//! `$OUT_DIR/mining.fatbin`. The driver picks the matching cubin at load;
|
||||
//! an unknown future card JIT-compiles the PTX.
|
||||
//! 3. If `nvcc` is not available (the `rust` lint runner, a workstation without
|
||||
//! CUDA), write an empty fat binary and warn. The crate still compiles;
|
||||
//! `CudaEngine::try_new` then fails with a clear message and the miner
|
||||
//! falls back to the wgpu engine.
|
||||
//!
|
||||
//! Knobs (environment):
|
||||
//! NVCC / CUDA_HOME where nvcc is
|
||||
//! MINER_CUDA_ARCHS comma-separated SM list, default 86,89,120
|
||||
//! MINER_NVCC_CCBIN host compiler for nvcc (-ccbin)
|
||||
//! MINER_CUDA_ALLOW_UNSUPPORTED_COMPILER=1 pass -allow-unsupported-compiler
|
||||
//! MINER_CUDA_REQUIRE=1 fail the build instead of warning when nvcc is missing
|
||||
|
||||
use std::env;
|
||||
use std::fs;
|
||||
use std::path::{Path, PathBuf};
|
||||
use std::process::Command;
|
||||
|
||||
fn write_constants(out: &Path) {
|
||||
use qp_poseidon_constants as c;
|
||||
let mut s = String::new();
|
||||
s.push_str("// Generated by engine-cuda/build.rs from qp-poseidon-constants. Do not edit.\n");
|
||||
s.push_str("#pragma once\n");
|
||||
let arr = |name: &str, v: &[u64]| {
|
||||
let body: Vec<String> = v.iter().map(|x| format!("0x{x:016x}ull")).collect();
|
||||
format!(
|
||||
"__device__ __constant__ unsigned long long {}[{}] = {{\n {}\n}};\n",
|
||||
name,
|
||||
v.len(),
|
||||
body.join(",\n ")
|
||||
)
|
||||
};
|
||||
let arr2 = |name: &str, v: &[[u64; 12]]| {
|
||||
let rows: Vec<String> = v
|
||||
.iter()
|
||||
.map(|r| {
|
||||
let body: Vec<String> = r.iter().map(|x| format!("0x{x:016x}ull")).collect();
|
||||
format!(" {{{}}}", body.join(", "))
|
||||
})
|
||||
.collect();
|
||||
format!(
|
||||
"__device__ __constant__ unsigned long long {}[{}][12] = {{\n{}\n}};\n",
|
||||
name,
|
||||
v.len(),
|
||||
rows.join(",\n")
|
||||
)
|
||||
};
|
||||
s.push_str(&arr("RC_INTERNAL", &c::POSEIDON2_INTERNAL_CONSTANTS_RAW));
|
||||
s.push_str(&arr2(
|
||||
"RC_INITIAL",
|
||||
&c::POSEIDON2_INITIAL_EXTERNAL_CONSTANTS_RAW,
|
||||
));
|
||||
s.push_str(&arr2(
|
||||
"RC_TERMINAL",
|
||||
&c::POSEIDON2_TERMINAL_EXTERNAL_CONSTANTS_RAW,
|
||||
));
|
||||
s.push_str(&arr("MDS_DIAG", &c::POSEIDON2_MATRIX_DIAG_12_RAW));
|
||||
s.push_str(&format!(
|
||||
"#define SPONGE_WIDTH {}\n#define N_EXTERNAL_HALF {}\n#define N_INTERNAL {}\n",
|
||||
c::SPONGE_WIDTH,
|
||||
c::POSEIDON2_EXTERNAL_ROUNDS / 2,
|
||||
c::POSEIDON2_INTERNAL_ROUNDS
|
||||
));
|
||||
fs::write(out.join("poseidon2_constants.cuh"), s).expect("write constants header");
|
||||
}
|
||||
|
||||
fn find_nvcc() -> Option<PathBuf> {
|
||||
if let Ok(p) = env::var("NVCC") {
|
||||
let p = PathBuf::from(p);
|
||||
if p.is_file() {
|
||||
return Some(p);
|
||||
}
|
||||
}
|
||||
if let Ok(home) = env::var("CUDA_HOME").or_else(|_| env::var("CUDA_PATH")) {
|
||||
let p = PathBuf::from(home).join("bin").join("nvcc");
|
||||
if p.is_file() {
|
||||
return Some(p);
|
||||
}
|
||||
}
|
||||
if let Ok(path) = env::var("PATH") {
|
||||
for dir in env::split_paths(&path) {
|
||||
let p = dir.join("nvcc");
|
||||
if p.is_file() {
|
||||
return Some(p);
|
||||
}
|
||||
}
|
||||
}
|
||||
// The toolkit the fleet drivers support; prefer an exact 13.0 over "latest".
|
||||
for cand in [
|
||||
"/usr/local/cuda-13.0",
|
||||
"/usr/local/cuda-13",
|
||||
"/usr/local/cuda",
|
||||
] {
|
||||
let p = PathBuf::from(cand).join("bin").join("nvcc");
|
||||
if p.is_file() {
|
||||
return Some(p);
|
||||
}
|
||||
}
|
||||
None
|
||||
}
|
||||
|
||||
fn main() {
|
||||
println!("cargo:rerun-if-changed=build.rs");
|
||||
println!("cargo:rerun-if-changed=src/kernels/mining.cu");
|
||||
for v in [
|
||||
"NVCC",
|
||||
"CUDA_HOME",
|
||||
"CUDA_PATH",
|
||||
"MINER_CUDA_ARCHS",
|
||||
"MINER_NVCC_CCBIN",
|
||||
"MINER_CUDA_ALLOW_UNSUPPORTED_COMPILER",
|
||||
"MINER_CUDA_REQUIRE",
|
||||
] {
|
||||
println!("cargo:rerun-if-env-changed={v}");
|
||||
}
|
||||
|
||||
let out = PathBuf::from(env::var("OUT_DIR").expect("OUT_DIR"));
|
||||
write_constants(&out);
|
||||
|
||||
let archs = env::var("MINER_CUDA_ARCHS").unwrap_or_else(|_| "86,89,120".to_string());
|
||||
let sms: Vec<String> = archs
|
||||
.split(',')
|
||||
.map(|s| s.trim().trim_start_matches("sm_").to_string())
|
||||
.filter(|s| !s.is_empty())
|
||||
.collect();
|
||||
println!("cargo:rustc-env=MINER_CUDA_ARCHS={}", sms.join(","));
|
||||
|
||||
let fatbin = out.join("mining.fatbin");
|
||||
let Some(nvcc) = find_nvcc() else {
|
||||
if env::var("MINER_CUDA_REQUIRE")
|
||||
.map(|v| v == "1")
|
||||
.unwrap_or(false)
|
||||
{
|
||||
panic!("nvcc not found and MINER_CUDA_REQUIRE=1; set NVCC or CUDA_HOME");
|
||||
}
|
||||
println!("cargo:warning=engine-cuda: nvcc not found; CUDA kernel not compiled into this binary (wgpu engine will be used)");
|
||||
fs::write(&fatbin, []).expect("write empty fatbin");
|
||||
return;
|
||||
};
|
||||
|
||||
let mut cmd = Command::new(&nvcc);
|
||||
cmd.arg("-fatbin")
|
||||
.arg("-O3")
|
||||
.arg("-std=c++17")
|
||||
.arg("-I")
|
||||
.arg(&out)
|
||||
.arg("-o")
|
||||
.arg(&fatbin);
|
||||
if let Ok(cc) = env::var("MINER_NVCC_CCBIN") {
|
||||
cmd.arg("-ccbin").arg(cc);
|
||||
}
|
||||
if env::var("MINER_CUDA_ALLOW_UNSUPPORTED_COMPILER")
|
||||
.map(|v| v == "1")
|
||||
.unwrap_or(false)
|
||||
{
|
||||
cmd.arg("-allow-unsupported-compiler");
|
||||
}
|
||||
for sm in &sms {
|
||||
cmd.arg("--generate-code")
|
||||
.arg(format!("arch=compute_{sm},code=sm_{sm}"));
|
||||
}
|
||||
if let Some(newest) = sms.iter().max_by_key(|s| s.parse::<u32>().unwrap_or(0)) {
|
||||
cmd.arg("--generate-code")
|
||||
.arg(format!("arch=compute_{newest},code=compute_{newest}"));
|
||||
}
|
||||
cmd.arg("src/kernels/mining.cu");
|
||||
|
||||
println!(
|
||||
"cargo:warning=engine-cuda: {} for sm_{{{}}}",
|
||||
nvcc.display(),
|
||||
sms.join(",")
|
||||
);
|
||||
let status = cmd.status().expect("run nvcc");
|
||||
if !status.success() {
|
||||
panic!("nvcc failed ({status}); see output above");
|
||||
}
|
||||
let size = fs::metadata(&fatbin).map(|m| m.len()).unwrap_or(0);
|
||||
if size == 0 {
|
||||
panic!("nvcc produced an empty fat binary");
|
||||
}
|
||||
}
|
||||
246
crates/engine-cuda/src/kernels/mining.cu
Normal file
246
crates/engine-cuda/src/kernels/mining.cu
Normal file
@@ -0,0 +1,246 @@
|
||||
// lair: Poseidon2-over-Goldilocks mining kernel, native CUDA (quantus/miner#3).
|
||||
//
|
||||
// Same host contract as engine-gpu's mining_u64.wgsl, and bit-exact with it and
|
||||
// with pow_core:
|
||||
// - the host precomputes the sponge midstate after absorbing the 32-byte
|
||||
// header and the high 32 bytes of the big-endian nonce (pow_core::
|
||||
// mining_midstate), so each nonce costs 2 permutations instead of 5;
|
||||
// - a batch never carries into the high 256 bits of the nonce, so only the
|
||||
// low 8 u32 limbs are incremented here;
|
||||
// - the first squeeze yields the most significant 256 bits of the hash, which
|
||||
// decide hash-vs-target on their own unless they exactly equal the target's
|
||||
// high half; only such candidates pay for the second squeeze.
|
||||
//
|
||||
// What CUDA buys over WGSL: a 64x64 -> 128-bit multiply is `__umul64hi` plus a
|
||||
// plain multiply, instead of four 32-bit partial products with carry
|
||||
// reconstruction. The field multiply is the whole cost of this hash.
|
||||
//
|
||||
// Field elements are kept in lazy (non-canonical) form below 2^64, exactly as
|
||||
// the WGSL kernel does, and canonicalised only when bytes are produced.
|
||||
|
||||
#include <stdint.h>
|
||||
#include "poseidon2_constants.cuh"
|
||||
|
||||
typedef unsigned long long u64;
|
||||
typedef unsigned int u32;
|
||||
|
||||
#define P64 0xFFFFFFFF00000001ull
|
||||
// 2^64 mod P = 2^32 - 1
|
||||
#define EPS 0xFFFFFFFFull
|
||||
|
||||
// a + b mod P, lazy. A wrapped carry folds back as 2^64 = EPS (mod P); the
|
||||
// second fold can only be needed when the first one wrapped.
|
||||
__device__ __forceinline__ u64 gf_add(u64 a, u64 b) {
|
||||
u64 s0 = a + b;
|
||||
u64 c1 = s0 < a;
|
||||
u64 s1 = s0 + (c1 ? EPS : 0ull);
|
||||
u64 c2 = c1 & (s1 < s0);
|
||||
return s1 + (c2 ? EPS : 0ull);
|
||||
}
|
||||
|
||||
// Reduce lo + hi * 2^64 mod P using 2^64 = EPS and 2^96 = -1 (mod P).
|
||||
__device__ __forceinline__ u64 gf_reduce(u64 lo, u64 hi) {
|
||||
u64 hi_hi = hi >> 32;
|
||||
u64 hi_lo = hi & EPS;
|
||||
u64 t0 = lo - hi_hi;
|
||||
if (lo < hi_hi) t0 -= EPS;
|
||||
u64 t1 = hi_lo * EPS;
|
||||
u64 t2 = t0 + t1;
|
||||
if (t2 < t0) t2 += EPS;
|
||||
return t2;
|
||||
}
|
||||
|
||||
__device__ __forceinline__ u64 gf_mul(u64 a, u64 b) {
|
||||
return gf_reduce(a * b, __umul64hi(a, b));
|
||||
}
|
||||
|
||||
__device__ __forceinline__ u64 gf_sqr(u64 a) {
|
||||
return gf_reduce(a * a, __umul64hi(a, a));
|
||||
}
|
||||
|
||||
// x^7
|
||||
__device__ __forceinline__ u64 gf_sbox(u64 x) {
|
||||
u64 x2 = gf_sqr(x);
|
||||
u64 x4 = gf_sqr(x2);
|
||||
u64 x6 = gf_mul(x4, x2);
|
||||
return gf_mul(x6, x);
|
||||
}
|
||||
|
||||
__device__ __forceinline__ u64 gf_canon(u64 a) {
|
||||
return a - (a >= P64 ? P64 : 0ull);
|
||||
}
|
||||
|
||||
// External linear layer: 4x4 MDS on each chunk, then circulant sums.
|
||||
__device__ __forceinline__ void ext_layer(u64 st[12]) {
|
||||
#pragma unroll
|
||||
for (int chunk = 0; chunk < 3; chunk++) {
|
||||
int o = chunk * 4;
|
||||
u64 x0 = st[o], x1 = st[o + 1], x2 = st[o + 2], x3 = st[o + 3];
|
||||
u64 t01 = gf_add(x0, x1);
|
||||
u64 t23 = gf_add(x2, x3);
|
||||
u64 t0123 = gf_add(t01, t23);
|
||||
u64 t01123 = gf_add(t0123, x1);
|
||||
u64 t01233 = gf_add(t0123, x3);
|
||||
st[o + 3] = gf_add(t01233, gf_add(x0, x0));
|
||||
st[o + 1] = gf_add(t01123, gf_add(x2, x2));
|
||||
st[o] = gf_add(t01123, t01);
|
||||
st[o + 2] = gf_add(t01233, t23);
|
||||
}
|
||||
u64 sums[4];
|
||||
#pragma unroll
|
||||
for (int k = 0; k < 4; k++) {
|
||||
sums[k] = gf_add(gf_add(st[k], st[k + 4]), st[k + 8]);
|
||||
}
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 12; i++) {
|
||||
st[i] = gf_add(st[i], sums[i & 3]);
|
||||
}
|
||||
}
|
||||
|
||||
// Internal linear layer: diagonal matrix plus full sum.
|
||||
__device__ __forceinline__ void int_layer(u64 st[12]) {
|
||||
u64 sum = st[0];
|
||||
#pragma unroll
|
||||
for (int i = 1; i < 12; i++) sum = gf_add(sum, st[i]);
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 12; i++) st[i] = gf_add(gf_mul(st[i], MDS_DIAG[i]), sum);
|
||||
}
|
||||
|
||||
__device__ __forceinline__ void permute(u64 st[12]) {
|
||||
ext_layer(st);
|
||||
#pragma unroll
|
||||
for (int r = 0; r < N_EXTERNAL_HALF; r++) {
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 12; i++) st[i] = gf_sbox(gf_add(st[i], RC_INITIAL[r][i]));
|
||||
ext_layer(st);
|
||||
}
|
||||
#pragma unroll 2
|
||||
for (int r = 0; r < N_INTERNAL; r++) {
|
||||
st[0] = gf_sbox(gf_add(st[0], RC_INTERNAL[r]));
|
||||
int_layer(st);
|
||||
}
|
||||
#pragma unroll
|
||||
for (int r = 0; r < N_EXTERNAL_HALF; r++) {
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 12; i++) st[i] = gf_sbox(gf_add(st[i], RC_TERMINAL[r][i]));
|
||||
ext_layer(st);
|
||||
}
|
||||
}
|
||||
|
||||
__device__ __forceinline__ u32 bswap32(u32 v) {
|
||||
return __byte_perm(v, 0, 0x0123);
|
||||
}
|
||||
|
||||
// Bindings mirror the WGSL kernel:
|
||||
// results [flag, nonce(16 u32 LE), hash(16 u32 LE)] = 33 u32
|
||||
// midstate 12 felts as u64
|
||||
// start_nonce 16 u32 LE limbs
|
||||
// target 16 u32 LE limbs
|
||||
extern "C" __global__ void __launch_bounds__(256)
|
||||
mining_main(u32* __restrict__ results,
|
||||
const u64* __restrict__ midstate,
|
||||
const u32* __restrict__ start_nonce,
|
||||
const u32* __restrict__ target,
|
||||
u32 total_threads,
|
||||
u32 nonces_per_thread,
|
||||
u32 total_nonces)
|
||||
{
|
||||
if (*((volatile u32*)results) != 0u) return;
|
||||
u32 tid = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (tid >= total_threads) return;
|
||||
u32 base_index = tid * nonces_per_thread;
|
||||
|
||||
u64 mid[12];
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 12; i++) mid[i] = midstate[i];
|
||||
u32 tgt[16];
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 16; i++) tgt[i] = target[i];
|
||||
u32 nonce_base[16];
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 16; i++) nonce_base[i] = start_nonce[i];
|
||||
|
||||
for (u32 j = 0; j < nonces_per_thread; j++) {
|
||||
u32 logical_index = base_index + j;
|
||||
if (logical_index >= total_nonces) break;
|
||||
if (j > 0u && *((volatile u32*)results) != 0u) return;
|
||||
|
||||
// Low 256 bits only; the host guarantees no carry into limbs 8..15.
|
||||
u32 current_nonce[16];
|
||||
u32 val0 = nonce_base[0];
|
||||
u32 sum0 = val0 + logical_index;
|
||||
current_nonce[0] = sum0;
|
||||
u32 carry = sum0 < val0 ? 1u : 0u;
|
||||
#pragma unroll
|
||||
for (int i = 1; i < 8; i++) {
|
||||
u32 val = nonce_base[i];
|
||||
u32 sum = val + carry;
|
||||
current_nonce[i] = sum;
|
||||
carry = sum < val ? 1u : 0u;
|
||||
}
|
||||
#pragma unroll
|
||||
for (int i = 8; i < 16; i++) current_nonce[i] = nonce_base[i];
|
||||
|
||||
// Resume the sponge from the midstate: absorb the low nonce half,
|
||||
// pad, squeeze twice (3 permutations instead of 5; the second squeeze
|
||||
// only for candidates).
|
||||
u64 st[12];
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 12; i++) st[i] = mid[i];
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 8; i++) st[i] = gf_add(st[i], (u64)bswap32(current_nonce[7 - i]));
|
||||
permute(st);
|
||||
st[0] = gf_add(st[0], 1ull);
|
||||
st[1] = gf_add(st[1], 1ull);
|
||||
permute(st);
|
||||
|
||||
// First squeeze: most significant 256 bits of the hash.
|
||||
u32 first[8];
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 4; i++) {
|
||||
u64 c = gf_canon(st[i]);
|
||||
first[2 * i] = (u32)(c & EPS);
|
||||
first[2 * i + 1] = (u32)(c >> 32);
|
||||
}
|
||||
u32 cmp = 0u; // 0 equal so far, 1 above target, 2 below target
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 8; i++) {
|
||||
u32 h = bswap32(first[i]);
|
||||
u32 t = tgt[15 - i];
|
||||
if (h != t) { cmp = h > t ? 1u : 2u; break; }
|
||||
}
|
||||
if (cmp == 1u) continue;
|
||||
|
||||
u32 hash_le[16];
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 8; i++) hash_le[15 - i] = bswap32(first[i]);
|
||||
permute(st);
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 4; i++) {
|
||||
u64 c = gf_canon(st[i]);
|
||||
hash_le[7 - 2 * i] = bswap32((u32)(c & EPS));
|
||||
hash_le[6 - 2 * i] = bswap32((u32)(c >> 32));
|
||||
}
|
||||
bool below = (cmp == 2u);
|
||||
if (!below) {
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 8; i++) {
|
||||
u32 h = hash_le[7 - i];
|
||||
u32 t = tgt[7 - i];
|
||||
if (h != t) { below = h < t; break; }
|
||||
}
|
||||
}
|
||||
|
||||
if (below) {
|
||||
if (atomicExch(&results[0], 1u) == 0u) {
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 16; i++) {
|
||||
results[1 + i] = current_nonce[i];
|
||||
results[17 + i] = hash_le[i];
|
||||
}
|
||||
__threadfence();
|
||||
}
|
||||
return;
|
||||
}
|
||||
}
|
||||
}
|
||||
487
crates/engine-cuda/src/lib.rs
Normal file
487
crates/engine-cuda/src/lib.rs
Normal file
@@ -0,0 +1,487 @@
|
||||
#![deny(rust_2018_idioms)]
|
||||
|
||||
//! lair: native CUDA mining engine behind `MinerEngine` (quantus/miner#3).
|
||||
//!
|
||||
//! Same host contract and batch loop as `engine-gpu` (midstate per batch, no
|
||||
//! carry into the high nonce half, cancellation between batches, thread-local
|
||||
//! worker-to-device assignment), with the kernel in `kernels/mining.cu` built
|
||||
//! into a fat binary by `build.rs`. Nothing in `engine-gpu` is touched; the
|
||||
//! service picks this engine when the binary carries a kernel and a CUDA
|
||||
//! driver is present, and falls back to wgpu otherwise.
|
||||
|
||||
use cudarc::driver::{
|
||||
CudaContext, CudaFunction, CudaSlice, CudaStream, DriverError, LaunchConfig, PushKernelArg,
|
||||
};
|
||||
use cudarc::nvrtc::Ptx;
|
||||
use engine_cpu::{CancelCheck, Candidate, EngineStatus, FoundOrigin, MinerEngine, Range};
|
||||
use pow_core::{format_hashrate, format_u512, JobContext};
|
||||
use primitive_types::U512;
|
||||
use std::cell::RefCell;
|
||||
use std::sync::atomic::{AtomicUsize, Ordering};
|
||||
use std::sync::Arc;
|
||||
use std::time::Instant;
|
||||
|
||||
/// The fat binary produced by build.rs; empty when nvcc was not available.
|
||||
const FATBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/mining.fatbin"));
|
||||
/// SM list the fat binary carries, for `--version` and the config metric.
|
||||
pub const CUDA_ARCHS: &str = env!("MINER_CUDA_ARCHS");
|
||||
const KERNEL_ID: &str = "cuda";
|
||||
const THREADS_PER_BLOCK: u32 = 256;
|
||||
/// Threads per SM the grid is sized for (override: MINER_CUDA_THREADS_PER_SM).
|
||||
/// Sets how many nonces each thread loops over for a given batch; tuned with #2.
|
||||
const THREADS_PER_SM_DEFAULT: u32 = 8192;
|
||||
|
||||
fn threads_per_sm() -> u32 {
|
||||
std::env::var("MINER_CUDA_THREADS_PER_SM")
|
||||
.ok()
|
||||
.and_then(|v| v.parse().ok())
|
||||
.filter(|v: &u32| *v > 0)
|
||||
.unwrap_or(THREADS_PER_SM_DEFAULT)
|
||||
}
|
||||
|
||||
static ENGINE_ID_COUNTER: AtomicUsize = AtomicUsize::new(0);
|
||||
|
||||
thread_local! {
|
||||
static ASSIGNED_DEVICE: RefCell<Option<(usize, usize)>> = const { RefCell::new(None) };
|
||||
static WORKER_RESOURCES: RefCell<Option<(usize, WorkerResources)>> = const { RefCell::new(None) };
|
||||
static DEVICE_LOST: RefCell<Option<usize>> = const { RefCell::new(None) };
|
||||
}
|
||||
|
||||
struct Device {
|
||||
ctx: Arc<CudaContext>,
|
||||
func: CudaFunction,
|
||||
name: String,
|
||||
sm_count: u32,
|
||||
threads_per_sm: u32,
|
||||
metrics: metrics::DeviceMetrics,
|
||||
}
|
||||
|
||||
struct WorkerResources {
|
||||
stream: Arc<CudaStream>,
|
||||
results: CudaSlice<u32>,
|
||||
midstate: CudaSlice<u64>,
|
||||
start_nonce: CudaSlice<u32>,
|
||||
target: CudaSlice<u32>,
|
||||
host_results: Vec<u32>,
|
||||
/// Target written for this job context; rewritten when it changes.
|
||||
target_written: Option<U512>,
|
||||
}
|
||||
|
||||
const RESULTS_LEN: usize = 1 + 16 + 16;
|
||||
|
||||
pub struct CudaEngine {
|
||||
engine_id: usize,
|
||||
devices: Vec<Arc<Device>>,
|
||||
device_counter: AtomicUsize,
|
||||
batch_size: u32,
|
||||
throttle_ms: u64,
|
||||
}
|
||||
|
||||
impl CudaEngine {
|
||||
/// Returns an error when the binary carries no kernel, no CUDA driver is
|
||||
/// present, or no device initialises.
|
||||
pub fn try_new(batch_size: u32, throttle_ms: u64) -> Result<Self, String> {
|
||||
if batch_size == 0 {
|
||||
return Err("batch size must be non-zero".into());
|
||||
}
|
||||
if FATBIN.is_empty() {
|
||||
return Err("CUDA kernel not compiled into this binary (built without nvcc)".into());
|
||||
}
|
||||
let count =
|
||||
CudaContext::device_count().map_err(|e| format!("CUDA driver unavailable: {e:?}"))?;
|
||||
if count <= 0 {
|
||||
return Err("no CUDA devices".into());
|
||||
}
|
||||
let mut devices = Vec::new();
|
||||
for ordinal in 0..count as usize {
|
||||
match Device::init(ordinal) {
|
||||
Ok(d) => {
|
||||
log::info!(
|
||||
target: "cuda_engine",
|
||||
"CUDA device {ordinal}: {} ({} SMs) using {KERNEL_ID} kernel [sm_{{{}}}]",
|
||||
d.name,
|
||||
d.sm_count,
|
||||
CUDA_ARCHS
|
||||
);
|
||||
devices.push(Arc::new(d));
|
||||
}
|
||||
Err(e) => {
|
||||
log::warn!(target: "cuda_engine", "CUDA device {ordinal} failed to initialise: {e:?}; skipping");
|
||||
}
|
||||
}
|
||||
}
|
||||
if devices.is_empty() {
|
||||
return Err("no CUDA device could be initialised".into());
|
||||
}
|
||||
log::info!(
|
||||
target: "cuda_engine",
|
||||
"CUDA engine initialized with {} devices (batch size: {} nonces, throttle: {}ms)",
|
||||
devices.len(),
|
||||
batch_size,
|
||||
throttle_ms
|
||||
);
|
||||
Ok(Self {
|
||||
engine_id: ENGINE_ID_COUNTER.fetch_add(1, Ordering::SeqCst),
|
||||
devices,
|
||||
device_counter: AtomicUsize::new(0),
|
||||
batch_size,
|
||||
throttle_ms,
|
||||
})
|
||||
}
|
||||
|
||||
pub fn device_count(&self) -> usize {
|
||||
self.devices.len()
|
||||
}
|
||||
|
||||
/// Drop the calling thread's device resources (call on worker exit).
|
||||
pub fn clear_worker_resources() {
|
||||
WORKER_RESOURCES.with(|r| *r.borrow_mut() = None);
|
||||
ASSIGNED_DEVICE.with(|a| *a.borrow_mut() = None);
|
||||
}
|
||||
}
|
||||
|
||||
impl Device {
|
||||
fn init(ordinal: usize) -> Result<Self, DriverError> {
|
||||
let ctx = CudaContext::new(ordinal)?;
|
||||
let name = ctx.name()?;
|
||||
let sm_count = ctx.attribute(
|
||||
cudarc::driver::sys::CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT,
|
||||
)? as u32;
|
||||
let module = ctx.load_module(Ptx::from_binary(FATBIN.to_vec()))?;
|
||||
let func = module.load_function("mining_main")?;
|
||||
Ok(Self {
|
||||
ctx,
|
||||
func,
|
||||
name,
|
||||
sm_count,
|
||||
threads_per_sm: threads_per_sm(),
|
||||
metrics: metrics::DeviceMetrics::new(ordinal, KERNEL_ID),
|
||||
})
|
||||
}
|
||||
|
||||
fn create_resources(&self) -> Result<WorkerResources, DriverError> {
|
||||
let stream = self.ctx.new_stream()?;
|
||||
Ok(WorkerResources {
|
||||
results: stream.alloc_zeros::<u32>(RESULTS_LEN)?,
|
||||
midstate: stream.alloc_zeros::<u64>(12)?,
|
||||
start_nonce: stream.alloc_zeros::<u32>(16)?,
|
||||
target: stream.alloc_zeros::<u32>(16)?,
|
||||
host_results: vec![0u32; RESULTS_LEN],
|
||||
target_written: None,
|
||||
stream,
|
||||
})
|
||||
}
|
||||
}
|
||||
|
||||
enum BatchResult {
|
||||
Found {
|
||||
candidate: Candidate,
|
||||
hash_count: u64,
|
||||
},
|
||||
NotFound {
|
||||
hash_count: u64,
|
||||
},
|
||||
DeviceLost,
|
||||
}
|
||||
|
||||
impl MinerEngine for CudaEngine {
|
||||
fn name(&self) -> &'static str {
|
||||
"gpu-cuda"
|
||||
}
|
||||
|
||||
fn prepare_context(&self, header_hash: [u8; 32], difficulty: U512) -> JobContext {
|
||||
JobContext::new(header_hash, difficulty)
|
||||
}
|
||||
|
||||
fn as_any(&self) -> &dyn std::any::Any {
|
||||
self
|
||||
}
|
||||
|
||||
fn search_range(
|
||||
&self,
|
||||
ctx: &JobContext,
|
||||
range: Range,
|
||||
cancel: &dyn CancelCheck,
|
||||
) -> EngineStatus {
|
||||
if DEVICE_LOST.with(|l| *l.borrow() == Some(self.engine_id)) {
|
||||
return EngineStatus::DeviceLost { hash_count: 0 };
|
||||
}
|
||||
if range.start > range.end {
|
||||
return EngineStatus::Exhausted { hash_count: 0 };
|
||||
}
|
||||
if cancel.is_cancelled() {
|
||||
return EngineStatus::Cancelled { hash_count: 0 };
|
||||
}
|
||||
|
||||
let device_index = ASSIGNED_DEVICE.with(|a| {
|
||||
let mut a = a.borrow_mut();
|
||||
match *a {
|
||||
Some((id, idx)) if id == self.engine_id => idx,
|
||||
_ => {
|
||||
let idx = if self.devices.len() == 1 {
|
||||
0
|
||||
} else {
|
||||
self.device_counter.fetch_add(1, Ordering::SeqCst) % self.devices.len()
|
||||
};
|
||||
*a = Some((self.engine_id, idx));
|
||||
log::info!(target: "cuda_engine", "Worker thread assigned to CUDA device {idx} (of {} total devices)", self.devices.len());
|
||||
idx
|
||||
}
|
||||
}
|
||||
});
|
||||
let dev = &self.devices[device_index];
|
||||
|
||||
WORKER_RESOURCES.with(|cell| {
|
||||
let mut slot = cell.borrow_mut();
|
||||
let need_new = !matches!(&*slot, Some((id, _)) if *id == self.engine_id);
|
||||
if need_new {
|
||||
match dev.create_resources() {
|
||||
Ok(r) => *slot = Some((self.engine_id, r)),
|
||||
Err(e) => {
|
||||
log::error!(target: "cuda_engine", "CUDA device {device_index} resource allocation failed: {e:?}");
|
||||
DEVICE_LOST.with(|l| *l.borrow_mut() = Some(self.engine_id));
|
||||
return EngineStatus::DeviceLost { hash_count: 0 };
|
||||
}
|
||||
}
|
||||
}
|
||||
let (_, res) = slot.as_mut().expect("resources present");
|
||||
self.search_on(dev, device_index, res, ctx, range, cancel)
|
||||
})
|
||||
}
|
||||
}
|
||||
|
||||
impl CudaEngine {
|
||||
fn search_on(
|
||||
&self,
|
||||
dev: &Device,
|
||||
device_index: usize,
|
||||
res: &mut WorkerResources,
|
||||
ctx: &JobContext,
|
||||
range: Range,
|
||||
cancel: &dyn CancelCheck,
|
||||
) -> EngineStatus {
|
||||
// Target is per job; write it once per context.
|
||||
if res.target_written != Some(ctx.target) {
|
||||
let target_bytes = ctx.target.to_little_endian();
|
||||
let mut target_u32s = [0u32; 16];
|
||||
for i in 0..16 {
|
||||
target_u32s[i] =
|
||||
u32::from_le_bytes(target_bytes[i * 4..(i + 1) * 4].try_into().unwrap());
|
||||
}
|
||||
if res
|
||||
.stream
|
||||
.memcpy_htod(&target_u32s, &mut res.target)
|
||||
.is_err()
|
||||
{
|
||||
return self.device_lost(dev, 0);
|
||||
}
|
||||
res.target_written = Some(ctx.target);
|
||||
}
|
||||
|
||||
let search_start = Instant::now();
|
||||
let mut total_hashes: u64 = 0;
|
||||
let mut current_start = range.start;
|
||||
let mut batch_num = 0u64;
|
||||
let mut last_batch_hashes: u64 = 0;
|
||||
|
||||
log::info!(
|
||||
target: "cuda_engine",
|
||||
"CUDA {} search started: range {}..{}, batch size: {} nonces",
|
||||
device_index,
|
||||
format_u512(range.start),
|
||||
format_u512(range.end),
|
||||
self.batch_size
|
||||
);
|
||||
|
||||
while current_start <= range.end {
|
||||
if cancel.is_cancelled() {
|
||||
dev.metrics.record_stale_hashes(last_batch_hashes);
|
||||
let elapsed = search_start.elapsed();
|
||||
log::info!(
|
||||
target: "cuda_engine",
|
||||
"CUDA {} cancelled before batch {} ({} total hashes in {:.2}s, {})",
|
||||
device_index,
|
||||
batch_num,
|
||||
total_hashes,
|
||||
elapsed.as_secs_f64(),
|
||||
format_hashrate(total_hashes as f64 / elapsed.as_secs_f64())
|
||||
);
|
||||
return EngineStatus::Cancelled {
|
||||
hash_count: total_hashes,
|
||||
};
|
||||
}
|
||||
|
||||
// Clamp so nonce increments never carry into the high 256 bits,
|
||||
// which the midstate precompute relies on.
|
||||
let remaining = range
|
||||
.end
|
||||
.saturating_sub(current_start)
|
||||
.saturating_add(U512::one());
|
||||
let headroom =
|
||||
(U512::one() << 256) - (current_start & ((U512::one() << 256) - U512::one()));
|
||||
let cap = remaining.min(headroom);
|
||||
let this_batch: u32 = if cap > U512::from(self.batch_size) {
|
||||
self.batch_size
|
||||
} else {
|
||||
cap.low_u32()
|
||||
};
|
||||
|
||||
match self.run_single_batch(dev, res, ctx, current_start, this_batch) {
|
||||
BatchResult::Found {
|
||||
candidate,
|
||||
hash_count,
|
||||
} => {
|
||||
total_hashes += hash_count;
|
||||
return EngineStatus::Found {
|
||||
candidate,
|
||||
hash_count: total_hashes,
|
||||
origin: FoundOrigin::GpuG1,
|
||||
};
|
||||
}
|
||||
BatchResult::NotFound { hash_count } => {
|
||||
total_hashes += hash_count;
|
||||
last_batch_hashes = hash_count;
|
||||
}
|
||||
BatchResult::DeviceLost => return self.device_lost(dev, total_hashes),
|
||||
}
|
||||
|
||||
current_start = current_start.saturating_add(U512::from(this_batch));
|
||||
batch_num += 1;
|
||||
|
||||
if self.throttle_ms > 0 && current_start <= range.end {
|
||||
let step = std::time::Duration::from_millis((self.throttle_ms / 10).max(1));
|
||||
let mut remaining = std::time::Duration::from_millis(self.throttle_ms);
|
||||
while remaining > std::time::Duration::ZERO {
|
||||
if cancel.is_cancelled() {
|
||||
return EngineStatus::Cancelled {
|
||||
hash_count: total_hashes,
|
||||
};
|
||||
}
|
||||
let s = remaining.min(step);
|
||||
std::thread::sleep(s);
|
||||
remaining = remaining.saturating_sub(s);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
let elapsed = search_start.elapsed();
|
||||
log::info!(
|
||||
target: "cuda_engine",
|
||||
"CUDA {} search exhausted: {} hashes in {} batches ({:.2}s, {})",
|
||||
device_index,
|
||||
total_hashes,
|
||||
batch_num,
|
||||
elapsed.as_secs_f64(),
|
||||
format_hashrate(total_hashes as f64 / elapsed.as_secs_f64())
|
||||
);
|
||||
EngineStatus::Exhausted {
|
||||
hash_count: total_hashes,
|
||||
}
|
||||
}
|
||||
|
||||
fn device_lost(&self, dev: &Device, hash_count: u64) -> EngineStatus {
|
||||
dev.metrics.record_device_lost();
|
||||
DEVICE_LOST.with(|l| *l.borrow_mut() = Some(self.engine_id));
|
||||
log::error!(target: "cuda_engine", "CUDA device lost or unresponsive - stopping worker");
|
||||
EngineStatus::DeviceLost { hash_count }
|
||||
}
|
||||
|
||||
fn run_single_batch(
|
||||
&self,
|
||||
dev: &Device,
|
||||
res: &mut WorkerResources,
|
||||
ctx: &JobContext,
|
||||
batch_start: U512,
|
||||
batch_size: u32,
|
||||
) -> BatchResult {
|
||||
let batch_start_at = Instant::now();
|
||||
|
||||
// Grid: enough threads to fill the card several waves deep, then loop
|
||||
// the remainder per thread.
|
||||
let max_threads = (dev.sm_count * dev.threads_per_sm) as u64;
|
||||
let logical_threads = (batch_size as u64).min(max_threads).max(1);
|
||||
let num_blocks = ((logical_threads as u32).div_ceil(THREADS_PER_BLOCK)).max(1);
|
||||
let total_threads = (num_blocks * THREADS_PER_BLOCK) as u64;
|
||||
let nonces_per_thread = ((batch_size as u64).div_ceil(total_threads)).max(1) as u32;
|
||||
let total_threads_u32 = total_threads as u32;
|
||||
|
||||
let start_nonce_bytes = batch_start.to_little_endian();
|
||||
let mut start_u32s = [0u32; 16];
|
||||
for i in 0..16 {
|
||||
start_u32s[i] =
|
||||
u32::from_le_bytes(start_nonce_bytes[i * 4..(i + 1) * 4].try_into().unwrap());
|
||||
}
|
||||
let nonce_be = batch_start.to_big_endian();
|
||||
let midstate = pow_core::mining_midstate(ctx.header, nonce_be[..32].try_into().unwrap());
|
||||
|
||||
let stream = res.stream.clone();
|
||||
let r: Result<(), DriverError> = (|| {
|
||||
stream.memcpy_htod(&start_u32s, &mut res.start_nonce)?;
|
||||
stream.memcpy_htod(&midstate, &mut res.midstate)?;
|
||||
stream.memset_zeros(&mut res.results)?;
|
||||
let cfg = LaunchConfig {
|
||||
grid_dim: (num_blocks, 1, 1),
|
||||
block_dim: (THREADS_PER_BLOCK, 1, 1),
|
||||
shared_mem_bytes: 0,
|
||||
};
|
||||
let mut launch = stream.launch_builder(&dev.func);
|
||||
launch
|
||||
.arg(&mut res.results)
|
||||
.arg(&res.midstate)
|
||||
.arg(&res.start_nonce)
|
||||
.arg(&res.target)
|
||||
.arg(&total_threads_u32)
|
||||
.arg(&nonces_per_thread)
|
||||
.arg(&batch_size);
|
||||
// SAFETY: the kernel signature matches the argument list above and
|
||||
// every pointer is a device buffer of at least the size the kernel reads.
|
||||
unsafe { launch.launch(cfg) }?;
|
||||
Ok(())
|
||||
})();
|
||||
if let Err(e) = r {
|
||||
log::error!(target: "cuda_engine", "CUDA batch submit failed: {e:?}");
|
||||
return BatchResult::DeviceLost;
|
||||
}
|
||||
let submitted_at = Instant::now();
|
||||
|
||||
if let Err(e) = stream
|
||||
.memcpy_dtoh(&res.results, &mut res.host_results)
|
||||
.and_then(|_| stream.synchronize())
|
||||
{
|
||||
log::error!(target: "cuda_engine", "CUDA batch readback failed: {e:?}");
|
||||
return BatchResult::DeviceLost;
|
||||
}
|
||||
let gpu_time = submitted_at.elapsed();
|
||||
dev.metrics
|
||||
.observe_batch(gpu_time, batch_start_at.elapsed() - gpu_time);
|
||||
|
||||
let out = &res.host_results;
|
||||
let dispatched = (total_threads * nonces_per_thread as u64).min(batch_size as u64);
|
||||
if out[0] != 0 {
|
||||
let nonce = U512::from_little_endian(&bytemuck_cast(&out[1..17]));
|
||||
let hash = U512::from_little_endian(&bytemuck_cast(&out[17..33]));
|
||||
let work = nonce.to_big_endian();
|
||||
let hashes_computed = if nonce >= batch_start {
|
||||
let logical_index = (nonce - batch_start).as_u64();
|
||||
let winning_iteration = logical_index % (nonces_per_thread as u64);
|
||||
(total_threads * (winning_iteration + 1)).min(dispatched)
|
||||
} else {
|
||||
dispatched
|
||||
};
|
||||
dev.metrics.record_hashes(hashes_computed);
|
||||
dev.metrics.record_solution();
|
||||
return BatchResult::Found {
|
||||
candidate: Candidate { nonce, work, hash },
|
||||
hash_count: hashes_computed,
|
||||
};
|
||||
}
|
||||
dev.metrics.record_hashes(dispatched);
|
||||
BatchResult::NotFound {
|
||||
hash_count: dispatched,
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
fn bytemuck_cast(words: &[u32]) -> Vec<u8> {
|
||||
words.iter().flat_map(|w| w.to_le_bytes()).collect()
|
||||
}
|
||||
@@ -87,6 +87,11 @@ enum Command {
|
||||
#[arg(long = "allow-integrated", env = "MINER_ALLOW_INTEGRATED")]
|
||||
allow_integrated: bool,
|
||||
|
||||
/// lair: GPU engine: auto (CUDA when this binary carries a kernel and a
|
||||
/// driver is present, else wgpu), cuda, or wgpu (quantus/miner#3).
|
||||
#[arg(long = "gpu-engine", env = "MINER_GPU_ENGINE", default_value = "auto")]
|
||||
gpu_engine: String,
|
||||
|
||||
/// Enable verbose logging
|
||||
#[arg(short, long, env = "MINER_VERBOSE")]
|
||||
verbose: bool,
|
||||
@@ -118,6 +123,11 @@ enum Command {
|
||||
#[arg(long = "allow-integrated", env = "MINER_ALLOW_INTEGRATED")]
|
||||
allow_integrated: bool,
|
||||
|
||||
/// lair: GPU engine: auto (CUDA when this binary carries a kernel and a
|
||||
/// driver is present, else wgpu), cuda, or wgpu (quantus/miner#3).
|
||||
#[arg(long = "gpu-engine", env = "MINER_GPU_ENGINE", default_value = "auto")]
|
||||
gpu_engine: String,
|
||||
|
||||
/// Enable verbose logging
|
||||
#[arg(short, long, env = "MINER_VERBOSE")]
|
||||
verbose: bool,
|
||||
@@ -171,6 +181,7 @@ async fn main() {
|
||||
gpu_throttle_ms,
|
||||
metrics_port,
|
||||
allow_integrated,
|
||||
gpu_engine,
|
||||
verbose,
|
||||
} => {
|
||||
init_logger(verbose);
|
||||
@@ -218,6 +229,7 @@ async fn main() {
|
||||
cpu_batch_size,
|
||||
gpu_throttle_ms,
|
||||
allow_integrated,
|
||||
gpu_engine,
|
||||
};
|
||||
|
||||
if let Err(e) = run(config).await {
|
||||
@@ -233,6 +245,7 @@ async fn main() {
|
||||
cpu_batch_size,
|
||||
duration,
|
||||
allow_integrated,
|
||||
gpu_engine,
|
||||
verbose,
|
||||
} => {
|
||||
init_logger(verbose);
|
||||
@@ -243,6 +256,7 @@ async fn main() {
|
||||
cpu_batch_size,
|
||||
duration,
|
||||
allow_integrated,
|
||||
gpu_engine,
|
||||
)
|
||||
.await;
|
||||
}
|
||||
@@ -304,9 +318,9 @@ fn init_logger(verbose: bool) {
|
||||
if std::env::var("RUST_LOG").is_err() {
|
||||
// Filter out noisy wgpu/naga shader compilation logs
|
||||
let log_level = if verbose {
|
||||
"debug,miner=debug,gpu_engine=debug,engine_cpu=debug,wgpu=warn,wgpu_core=warn,wgpu_hal=warn,naga=warn"
|
||||
"debug,miner=debug,gpu_engine=debug,cuda_engine=debug,engine_cpu=debug,wgpu=warn,wgpu_core=warn,wgpu_hal=warn,naga=warn"
|
||||
} else {
|
||||
"info,miner=info,gpu_engine=info,wgpu=error,wgpu_core=error,wgpu_hal=error,naga=error"
|
||||
"info,miner=info,gpu_engine=info,cuda_engine=info,wgpu=error,wgpu_core=error,wgpu_hal=error,naga=error"
|
||||
};
|
||||
std::env::set_var("RUST_LOG", log_level);
|
||||
}
|
||||
@@ -320,6 +334,7 @@ async fn run_benchmark(
|
||||
cpu_batch_size: u64,
|
||||
duration: u64,
|
||||
allow_integrated: bool,
|
||||
gpu_engine_pref: String,
|
||||
) {
|
||||
let effective_cpu_workers = cpu_workers.unwrap_or_else(num_cpus::get);
|
||||
|
||||
@@ -329,6 +344,7 @@ async fn run_benchmark(
|
||||
gpu_batch_size,
|
||||
0,
|
||||
allow_integrated,
|
||||
&gpu_engine_pref,
|
||||
) {
|
||||
Ok((engine, count)) => (engine, count),
|
||||
Err(e) => {
|
||||
|
||||
@@ -31,4 +31,5 @@ quantus-miner-api = { workspace = true }
|
||||
pow-core = { path = "../pow-core" }
|
||||
engine-cpu = { path = "../engine-cpu", optional = true }
|
||||
engine-gpu = { path = "../engine-gpu" }
|
||||
engine-cuda = { path = "../engine-cuda" } # lair: quantus/miner#3
|
||||
metrics = { path = "../metrics" }
|
||||
|
||||
@@ -39,6 +39,9 @@ pub struct ServiceConfig {
|
||||
pub gpu_throttle_ms: u64,
|
||||
/// Allow integrated GPUs even when discrete GPUs are available
|
||||
pub allow_integrated: bool,
|
||||
/// lair: GPU engine preference: "auto" (CUDA when available, else wgpu),
|
||||
/// "cuda" or "wgpu" (quantus/miner#3).
|
||||
pub gpu_engine: String,
|
||||
}
|
||||
|
||||
/// Engine type for tracking metrics per compute type.
|
||||
@@ -477,6 +480,7 @@ fn worker_loop(
|
||||
// Clean up GPU resources on thread exit
|
||||
if engine_type == EngineType::Gpu {
|
||||
engine_gpu::GpuEngine::clear_worker_resources();
|
||||
engine_cuda::CudaEngine::clear_worker_resources(); // lair
|
||||
}
|
||||
|
||||
log::debug!("{type_str} worker {thread_id} exited");
|
||||
@@ -488,12 +492,37 @@ pub fn resolve_gpu_configuration(
|
||||
batch_size: u32,
|
||||
throttle_ms: u64,
|
||||
allow_integrated: bool,
|
||||
gpu_engine: &str,
|
||||
) -> anyhow::Result<(Option<Arc<dyn MinerEngine>>, usize)> {
|
||||
// Explicit 0 means no GPU
|
||||
if requested_devices == Some(0) {
|
||||
return Ok((None, 0));
|
||||
}
|
||||
|
||||
// lair: native CUDA engine first unless wgpu was asked for (quantus/miner#3).
|
||||
if gpu_engine != "wgpu" {
|
||||
match engine_cuda::CudaEngine::try_new(batch_size, throttle_ms) {
|
||||
Ok(engine) => {
|
||||
let available = engine.device_count();
|
||||
let count = match requested_devices {
|
||||
Some(n) if n > available => anyhow::bail!(
|
||||
"Requested {n} GPU devices but only {available} available (CUDA)"
|
||||
),
|
||||
Some(n) => n,
|
||||
None => {
|
||||
log::info!("Auto-detected {available} CUDA device(s)");
|
||||
available
|
||||
}
|
||||
};
|
||||
return Ok((Some(Arc::new(engine)), count));
|
||||
}
|
||||
Err(e) if gpu_engine == "cuda" => {
|
||||
anyhow::bail!("CUDA engine requested but unavailable: {e}")
|
||||
}
|
||||
Err(e) => log::info!("CUDA engine unavailable ({e}); using wgpu"),
|
||||
}
|
||||
}
|
||||
|
||||
// Try to initialize GPU engine
|
||||
let engine = engine_gpu::GpuEngine::try_new(batch_size, throttle_ms, allow_integrated);
|
||||
let engine = match engine {
|
||||
@@ -541,6 +570,7 @@ pub async fn run(config: ServiceConfig) -> anyhow::Result<()> {
|
||||
config.gpu_batch_size,
|
||||
config.gpu_throttle_ms,
|
||||
config.allow_integrated,
|
||||
&config.gpu_engine,
|
||||
)?;
|
||||
|
||||
// Resolve CPU workers
|
||||
|
||||
Reference in New Issue
Block a user