diff --git a/.gitea/workflows/deploy.yaml b/.gitea/workflows/deploy.yaml index a909925..e98a234 100644 --- a/.gitea/workflows/deploy.yaml +++ b/.gitea/workflows/deploy.yaml @@ -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 diff --git a/Cargo.lock b/Cargo.lock index 0537959..f3a501f 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -107,7 +107,7 @@ version = "0.38.0+1.3.281" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "0bb44936d800fea8f016d7f2311c6a4f97aebd5dc86f09906139ec848cf3a46f" dependencies = [ - "libloading", + "libloading 0.8.9", ] [[package]] @@ -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]] diff --git a/Cargo.toml b/Cargo.toml index 3f4d1df..a650c82 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -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", diff --git a/crates/bench-harness/Cargo.toml b/crates/bench-harness/Cargo.toml index 1c5461c..4b64df4 100644 --- a/crates/bench-harness/Cargo.toml +++ b/crates/bench-harness/Cargo.toml @@ -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"] } diff --git a/crates/bench-harness/src/main.rs b/crates/bench-harness/src/main.rs index fe14ab4..0657457 100644 --- a/crates/bench-harness/src/main.rs +++ b/crates/bench-harness/src/main.rs @@ -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, } +/// 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 { 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, diff --git a/crates/engine-cuda/Cargo.toml b/crates/engine-cuda/Cargo.toml new file mode 100644 index 0000000..806854e --- /dev/null +++ b/crates/engine-cuda/Cargo.toml @@ -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" diff --git a/crates/engine-cuda/build.rs b/crates/engine-cuda/build.rs new file mode 100644 index 0000000..c3d6337 --- /dev/null +++ b/crates/engine-cuda/build.rs @@ -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 = 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 = v + .iter() + .map(|r| { + let body: Vec = 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 { + 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 = 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::().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"); + } +} diff --git a/crates/engine-cuda/src/kernels/mining.cu b/crates/engine-cuda/src/kernels/mining.cu new file mode 100644 index 0000000..178f9a8 --- /dev/null +++ b/crates/engine-cuda/src/kernels/mining.cu @@ -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 +#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; + } + } +} diff --git a/crates/engine-cuda/src/lib.rs b/crates/engine-cuda/src/lib.rs new file mode 100644 index 0000000..b7b7aca --- /dev/null +++ b/crates/engine-cuda/src/lib.rs @@ -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> = const { RefCell::new(None) }; + static WORKER_RESOURCES: RefCell> = const { RefCell::new(None) }; + static DEVICE_LOST: RefCell> = const { RefCell::new(None) }; +} + +struct Device { + ctx: Arc, + func: CudaFunction, + name: String, + sm_count: u32, + threads_per_sm: u32, + metrics: metrics::DeviceMetrics, +} + +struct WorkerResources { + stream: Arc, + results: CudaSlice, + midstate: CudaSlice, + start_nonce: CudaSlice, + target: CudaSlice, + host_results: Vec, + /// Target written for this job context; rewritten when it changes. + target_written: Option, +} + +const RESULTS_LEN: usize = 1 + 16 + 16; + +pub struct CudaEngine { + engine_id: usize, + devices: Vec>, + 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 { + 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 { + 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 { + let stream = self.ctx.new_stream()?; + Ok(WorkerResources { + results: stream.alloc_zeros::(RESULTS_LEN)?, + midstate: stream.alloc_zeros::(12)?, + start_nonce: stream.alloc_zeros::(16)?, + target: stream.alloc_zeros::(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 { + words.iter().flat_map(|w| w.to_le_bytes()).collect() +} diff --git a/crates/miner-cli/src/main.rs b/crates/miner-cli/src/main.rs index dc59517..1017672 100644 --- a/crates/miner-cli/src/main.rs +++ b/crates/miner-cli/src/main.rs @@ -87,6 +87,11 @@ enum Command { #[arg(long = "allow-integrated", env = "MINER_ALLOW_INTEGRATED")] allow_integrated: bool, + /// lair: GPU engine: auto (CUDA when this binary carries a kernel and a + /// driver is present, else wgpu), cuda, or wgpu (quantus/miner#3). + #[arg(long = "gpu-engine", env = "MINER_GPU_ENGINE", default_value = "auto")] + gpu_engine: String, + /// Enable verbose logging #[arg(short, long, env = "MINER_VERBOSE")] verbose: bool, @@ -118,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) => { diff --git a/crates/miner-service/Cargo.toml b/crates/miner-service/Cargo.toml index 07a73c9..96d0be7 100644 --- a/crates/miner-service/Cargo.toml +++ b/crates/miner-service/Cargo.toml @@ -31,4 +31,5 @@ quantus-miner-api = { workspace = true } pow-core = { path = "../pow-core" } engine-cpu = { path = "../engine-cpu", optional = true } engine-gpu = { path = "../engine-gpu" } +engine-cuda = { path = "../engine-cuda" } # lair: quantus/miner#3 metrics = { path = "../metrics" } diff --git a/crates/miner-service/src/lib.rs b/crates/miner-service/src/lib.rs index c94677c..0ef9cfd 100644 --- a/crates/miner-service/src/lib.rs +++ b/crates/miner-service/src/lib.rs @@ -39,6 +39,9 @@ pub struct ServiceConfig { pub gpu_throttle_ms: u64, /// Allow integrated GPUs even when discrete GPUs are available pub allow_integrated: bool, + /// lair: GPU engine preference: "auto" (CUDA when available, else wgpu), + /// "cuda" or "wgpu" (quantus/miner#3). + pub gpu_engine: String, } /// Engine type for tracking metrics per compute type. @@ -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>, 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