diff --git a/.gitea/workflows/deploy.yaml b/.gitea/workflows/deploy.yaml index 87b5cd9..f40bfbe 100644 --- a/.gitea/workflows/deploy.yaml +++ b/.gitea/workflows/deploy.yaml @@ -88,17 +88,24 @@ jobs: # `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. + # `cuda_kernel` selects the kernel entry (MINER_CUDA_KERNEL, #29): + # unrolled is the default; loop measured +1.7% on the 4090 and + # -1.8% / -3.4% on the 3060 / 5090 (2026-09-14), so only benjy runs + # it. `kernel` must be the id that choice reports: cuda or cuda-loop. - host: benjy.hanzalova.internal node: bob.hanzalova.internal gpu_devices: "1" # 1x RTX 4090 (sm_89); reference card for #2 - kernel: cuda + cuda_kernel: loop + kernel: cuda-loop - host: quadbrat.hanzalova.internal node: bob.hanzalova.internal gpu_devices: "1" # 1x RTX 3060 (sm_86) + cuda_kernel: unrolled kernel: cuda - host: beast.hanzalova.internal node: bob.hanzalova.internal gpu_devices: "2" # 2x RTX 5090 (sm_120) + cuda_kernel: unrolled kernel: cuda steps: - uses: actions/checkout@v4 @@ -146,6 +153,7 @@ jobs: if: ${{ github.event.inputs.mode != 'validate' }} env: GPU_DEVICES: ${{ matrix.gpu_devices }} + CUDA_KERNEL: ${{ matrix.cuda_kernel }} run: | set -euo pipefail SSHOPTS="-i $HOME/.ssh/id_gitea_ci -o StrictHostKeyChecking=accept-new" @@ -210,6 +218,7 @@ jobs: t = pathlib.Path("deploy/miner.env.tmpl").read_text() t = t.replace("{{QUANTUS_NODE_ADDR}}", os.environ["NODE_ADDR"]) t = t.replace("{{QUANTUS_GPU_DEVICES}}", os.environ["GPU_DEVICES"]) + t = t.replace("{{MINER_CUDA_KERNEL}}", os.environ["CUDA_KERNEL"]) pathlib.Path("miner.env").write_text(t) PY push --chown=root:quantus-miner --chmod=F0640 \ diff --git a/crates/engine-cuda/src/kernels/mining.cu b/crates/engine-cuda/src/kernels/mining.cu index a2fc02d..f6e4ffd 100644 --- a/crates/engine-cuda/src/kernels/mining.cu +++ b/crates/engine-cuda/src/kernels/mining.cu @@ -333,6 +333,32 @@ __device__ __forceinline__ u64 gf_canon(u64 a) { #else #define LAIR_LAUNCH_BOUNDS __launch_bounds__(LAIR_TPB) #endif +// lair: the loop-structured kernel (`mining_loop`, quantus/miner#27). Same +// arithmetic and same host contract as `mining_main`, but the round loops +// are rolled with the constants fetched from constant memory by index, the +// way the fastest closed kernel is built, so the register footprint can drop +// and more warps fit per SM. Selected at run time (MINER_CUDA_KERNEL=loop); +// the unrolled kernel stays the default until this one measures better. +// 512 measured best for this kernel on beast (2131 MH/s against 2095 at +// 256 and 2099 at 128 x 4 blocks); the host picks the same block size when +// the loop kernel is selected. +#ifndef LAIR_LOOP_TPB +#define LAIR_LOOP_TPB 512 +#endif +#ifndef LAIR_LOOP_MINBLOCKS +#define LAIR_LOOP_MINBLOCKS 0 +#endif +#if LAIR_LOOP_MINBLOCKS +#define LAIR_LOOP_LAUNCH_BOUNDS __launch_bounds__(LAIR_LOOP_TPB, LAIR_LOOP_MINBLOCKS) +#else +#define LAIR_LOOP_LAUNCH_BOUNDS __launch_bounds__(LAIR_LOOP_TPB) +#endif +// Internal rounds per loop iteration in the loop kernel (1 or 2). One +// measured 2105 MH/s against 1976 for two on beast: the pair doubles the +// live state across the iteration for no scheduling gain. +#ifndef LAIR_LOOP_INT_PAIR +#define LAIR_LOOP_INT_PAIR 1 +#endif // Decide most nonces on element 0 of the final state before computing the // rest of the last linear layer: at mainnet difficulty the top ~47 bits of // the hash must be zero, and they live in that element. @@ -708,6 +734,65 @@ __device__ __forceinline__ void permute(u64 st[12], bool skip_last_ext = false, #endif } +// lair: rolled-loop twin of permute for the loop kernel. Round-constant +// rows are addressed by the loop variable, so they are constant-memory loads +// with a uniform offset instead of immediates; the per-element loops inside +// each round stay unrolled so the state lives in registers. +#if LAIR_LAZY_ADD && LAIR_FOLD_RC +static_assert(N_INTERNAL % 2 == 0, "the paired internal loop assumes an even round count"); +__device__ __forceinline__ void permute_loop(u64 st[12], bool skip_last_ext, bool skip_first_ext) { + if (!skip_first_ext) ext_layer_rc(st, RC_INITIAL[0]); +#pragma unroll 1 + for (int r = 0; r + 1 < N_EXTERNAL_HALF; r++) { +#pragma unroll + for (int i = 0; i < 12; i++) st[i] = gf_sbox(st[i]); + ext_layer_rc(st, RC_INITIAL[r + 1]); + } +#pragma unroll + for (int i = 0; i < 12; i++) st[i] = gf_sbox(st[i]); + ext_layer(st); + st[0] = gf_add(st[0], RC_INTERNAL[0]); +#if LAIR_LOOP_INT_PAIR == 2 +#pragma unroll 1 + for (int r = 0; r + 2 < N_INTERNAL; r += 2) { + st[0] = gf_sbox(st[0]); + int_layer_rc(st, RC_INTERNAL[r + 1], true, nullptr); + st[0] = gf_sbox(st[0]); + int_layer_rc(st, RC_INTERNAL[r + 2], true, nullptr); + } + st[0] = gf_sbox(st[0]); + int_layer_rc(st, RC_INTERNAL[N_INTERNAL - 1], true, nullptr); +#else +#pragma unroll 1 + for (int r = 0; r + 1 < N_INTERNAL; r++) { + st[0] = gf_sbox(st[0]); + int_layer_rc(st, RC_INTERNAL[r + 1], true, nullptr); + } +#endif + st[0] = gf_sbox(st[0]); + int_layer_rc(st, 0ull, false, RC_TERMINAL[0]); +#pragma unroll 1 + for (int r = 0; r + 1 < N_EXTERNAL_HALF; r++) { +#pragma unroll + for (int i = 0; i < 12; i++) st[i] = gf_sbox(st[i]); + ext_layer_rc(st, RC_TERMINAL[r + 1]); + } +#pragma unroll + for (int i = 0; i < 12; i++) st[i] = gf_sbox(st[i]); + if (!skip_last_ext) ext_layer(st); +} +#else +__device__ __forceinline__ void permute_loop(u64 st[12], bool skip_last_ext, bool skip_first_ext) { + permute(st, skip_last_ext, skip_first_ext); +} +#endif + +template +__device__ __forceinline__ void run_permute(u64 st[12], bool skip_last_ext = false, bool skip_first_ext = false) { + if (LOOP) permute_loop(st, skip_last_ext, skip_first_ext); + else permute(st, skip_last_ext, skip_first_ext); +} + __device__ __forceinline__ u32 bswap32(u32 v) { return __byte_perm(v, 0, 0x0123); } @@ -742,12 +827,14 @@ static_assert(sizeof(MiningUniforms) == 320, "MiningUniforms must match the host __device__ __constant__ u64 LAIR_DIR7[12] = {1, 1, 3, 2, 2, 2, 6, 4, 1, 1, 3, 2}; #endif -extern "C" __global__ void LAIR_LAUNCH_BOUNDS -mining_main(u32* __restrict__ results, - const MiningUniforms uni, - u32 total_threads, - u32 nonces_per_thread, - u32 total_nonces) +// The per-nonce search, shared by both kernel entries; LOOP picks the +// permutation implementation at compile time. +template +__device__ __forceinline__ void mine_nonces(u32* __restrict__ results, + const MiningUniforms& uni, + u32 total_threads, + u32 nonces_per_thread, + u32 total_nonces) { if (*((volatile u32*)results) != 0u) return; u32 tid = blockIdx.x * blockDim.x + threadIdx.x; @@ -796,13 +883,13 @@ mining_main(u32* __restrict__ results, for (int i = 0; i < 8; i++) st[i] = gf_add(st[i], (u64)bswap32(current_nonce[7 - i])); ext_layer_rc(st, RC_INITIAL[0]); } - permute(st, false, true); + run_permute(st, false, true); #else #pragma unroll for (int i = 0; i < 12; i++) st[i] = uni.midstate[i]; #pragma unroll for (int i = 0; i < 8; i++) st[i] = gf_add(st[i], (u64)bswap32(current_nonce[7 - i])); - permute(st); + run_permute(st); #endif st[0] = gf_add(st[0], 1ull); st[1] = gf_add(st[1], 1ull); @@ -810,7 +897,7 @@ mining_main(u32* __restrict__ results, // Element 0 of the final state is the hash's top 64 bits. Compute it // alone, and only nonces that do not already exceed the target there // pay for the other eleven outputs of the last linear layer. - permute(st, true); + run_permute(st, true); { u64 c0 = gf_canon(ext_layer_out0(st)); u32 h0 = bswap32((u32)(c0 & EPS)); @@ -820,7 +907,7 @@ mining_main(u32* __restrict__ results, } ext_layer(st); #else - permute(st); + run_permute(st); #endif // First squeeze: most significant 256 bits of the hash. @@ -843,7 +930,7 @@ mining_main(u32* __restrict__ results, u32 hash_le[16]; #pragma unroll for (int i = 0; i < 8; i++) hash_le[15 - i] = bswap32(first[i]); - permute(st); + run_permute(st); #pragma unroll for (int i = 0; i < 4; i++) { u64 c = gf_canon(st[i]); @@ -874,3 +961,23 @@ mining_main(u32* __restrict__ results, } } } + +extern "C" __global__ void LAIR_LAUNCH_BOUNDS +mining_main(u32* __restrict__ results, + const MiningUniforms uni, + u32 total_threads, + u32 nonces_per_thread, + u32 total_nonces) +{ + mine_nonces(results, uni, total_threads, nonces_per_thread, total_nonces); +} + +extern "C" __global__ void LAIR_LOOP_LAUNCH_BOUNDS +mining_loop(u32* __restrict__ results, + const MiningUniforms uni, + u32 total_threads, + u32 nonces_per_thread, + u32 total_nonces) +{ + mine_nonces(results, uni, total_threads, nonces_per_thread, total_nonces); +} diff --git a/crates/engine-cuda/src/lib.rs b/crates/engine-cuda/src/lib.rs index 96aecd2..d6fa1a5 100644 --- a/crates/engine-cuda/src/lib.rs +++ b/crates/engine-cuda/src/lib.rs @@ -26,18 +26,60 @@ use std::time::Instant; 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"; +/// lair: which kernel entry the module runs (quantus/miner#27). `unrolled` +/// is `mining_main`, the deployed kernel; `loop` is `mining_loop`, the +/// rolled-round twin under evaluation. Selected with MINER_CUDA_KERNEL so a +/// host can be switched without a rebuild, and reported as the kernel id on +/// the per-device metrics so the deploy validate and the dashboards can tell +/// them apart. +#[derive(Clone, Copy, PartialEq, Eq, Debug)] +enum KernelChoice { + Unrolled, + Loop, +} + +impl KernelChoice { + fn from_env() -> Self { + match std::env::var("MINER_CUDA_KERNEL").as_deref() { + Ok("loop") => KernelChoice::Loop, + Ok("unrolled") | Ok("") | Err(_) => KernelChoice::Unrolled, + Ok(other) => { + log::warn!(target: "cuda_engine", "MINER_CUDA_KERNEL={other:?} is not unrolled|loop; using unrolled"); + KernelChoice::Unrolled + } + } + } + fn function(self) -> &'static str { + match self { + KernelChoice::Unrolled => "mining_main", + KernelChoice::Loop => "mining_loop", + } + } + /// Kernel id on metrics: `cuda` for the deployed kernel (what the deploy + /// validate asserts), `cuda-loop` for the candidate. + fn id(self) -> &'static str { + match self { + KernelChoice::Unrolled => "cuda", + KernelChoice::Loop => "cuda-loop", + } + } +} /// Threads per block. Must match the kernel's `__launch_bounds__` (LAIR_TPB, /// default 256); override for experiments with MINER_CUDA_THREADS_PER_BLOCK /// on a kernel compiled with the same value. const THREADS_PER_BLOCK_DEFAULT: u32 = 256; +/// The loop kernel's `LAIR_LOOP_TPB`. +const LOOP_THREADS_PER_BLOCK_DEFAULT: u32 = 512; -fn threads_per_block() -> u32 { +fn threads_per_block(kernel: KernelChoice) -> u32 { std::env::var("MINER_CUDA_THREADS_PER_BLOCK") .ok() .and_then(|v| v.parse().ok()) .filter(|v: &u32| *v > 0 && v.is_multiple_of(32)) - .unwrap_or(THREADS_PER_BLOCK_DEFAULT) + .unwrap_or(match kernel { + KernelChoice::Unrolled => THREADS_PER_BLOCK_DEFAULT, + KernelChoice::Loop => LOOP_THREADS_PER_BLOCK_DEFAULT, + }) } /// lair: size every batch to whole grids. The CLI default batch (1M nonces) @@ -93,6 +135,7 @@ struct Device { func: CudaFunction, name: String, sm_count: u32, + kernel: KernelChoice, threads_per_block: u32, threads_per_sm: u32, metrics: metrics::DeviceMetrics, @@ -171,9 +214,10 @@ impl CudaEngine { Ok(d) => { log::info!( target: "cuda_engine", - "CUDA device {ordinal}: {} ({} SMs) using {KERNEL_ID} kernel [sm_{{{}}}]", + "CUDA device {ordinal}: {} ({} SMs) using {} kernel [sm_{{{}}}]", d.name, d.sm_count, + d.kernel.id(), CUDA_ARCHS ); devices.push(Arc::new(d)); @@ -221,15 +265,17 @@ impl Device { 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")?; + let kernel = KernelChoice::from_env(); + let func = module.load_function(kernel.function())?; Ok(Self { ctx, func, name, sm_count, + kernel, threads_per_sm: threads_per_sm(), - threads_per_block: threads_per_block(), - metrics: metrics::DeviceMetrics::new(ordinal, KERNEL_ID), + threads_per_block: threads_per_block(kernel), + metrics: metrics::DeviceMetrics::new(ordinal, kernel.id()), }) } diff --git a/deploy/miner.env.tmpl b/deploy/miner.env.tmpl index efdb334..7b05468 100644 --- a/deploy/miner.env.tmpl +++ b/deploy/miner.env.tmpl @@ -4,3 +4,5 @@ # from the node host as separate 0640 files. QUANTUS_NODE_ADDR={{QUANTUS_NODE_ADDR}} QUANTUS_GPU_DEVICES={{QUANTUS_GPU_DEVICES}} +# Kernel entry in the CUDA engine: unrolled (mining_main) or loop (mining_loop), quantus/miner#29. +MINER_CUDA_KERNEL={{MINER_CUDA_KERNEL}}