diff --git a/Cargo.lock b/Cargo.lock index aa6200a..e618f9d 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -472,6 +472,7 @@ dependencies = [ "anyhow", "cust", "engine-cpu", + "hex", "log", "metrics", "num-bigint", diff --git a/crates/engine-gpu-cuda/Cargo.toml b/crates/engine-gpu-cuda/Cargo.toml index 20102da..48b2fc1 100644 --- a/crates/engine-gpu-cuda/Cargo.toml +++ b/crates/engine-gpu-cuda/Cargo.toml @@ -23,6 +23,7 @@ primitive-types = { workspace = true } log = { workspace = true } thiserror = { workspace = true } anyhow = { workspace = true } +hex = { workspace = true } # Metrics (optional; used for emitting generic counters/gauges from this engine) metrics = { path = "../metrics", optional = true } diff --git a/crates/engine-gpu-cuda/src/kernels/qpow_kernel.cu b/crates/engine-gpu-cuda/src/kernels/qpow_kernel.cu index 5a97c93..c9ee6e2 100644 --- a/crates/engine-gpu-cuda/src/kernels/qpow_kernel.cu +++ b/crates/engine-gpu-cuda/src/kernels/qpow_kernel.cu @@ -538,6 +538,9 @@ extern "C" __global__ void qpow_montgomery_g2_kernel( int* __restrict__ found_flag, // 0 -> not found, 1 -> found uint32_t* __restrict__ out_index, // linear index (t * iters + j) uint8_t* __restrict__ out_distance_be, // 64 bytes + // Debug output buffers (optional; host may pass nullptrs) + uint8_t* __restrict__ out_dbg_y_be, // 64 bytes (optional) + uint8_t* __restrict__ out_dbg_h_be, // 64 bytes (optional) // Threading parameters const uint32_t num_threads, @@ -648,15 +651,27 @@ extern "C" __global__ void qpow_montgomery_g2_kernel( if (out_index) { *out_index = tid * iters + j; } - // Write distance + // Write distance and debug buffers (if provided) if (out_distance_be) { -uint8_t dist_out[64]; + uint8_t dist_out[64]; be_u64x8_to_be64_bytes(dist_limb, dist_out); #pragma unroll for (int i = 0; i < 64; ++i) { out_distance_be[i] = dist_out[i]; } } + if (out_dbg_y_be) { +#pragma unroll + for (int i = 0; i < 64; ++i) { + out_dbg_y_be[i] = y_be[i]; + } + } + if (out_dbg_h_be) { +#pragma unroll + for (int i = 0; i < 64; ++i) { + out_dbg_h_be[i] = h_be[i]; + } + } } return; // early-exit after claiming } diff --git a/crates/engine-gpu-cuda/src/lib.rs b/crates/engine-gpu-cuda/src/lib.rs index 03d9588..5362c9c 100644 --- a/crates/engine-gpu-cuda/src/lib.rs +++ b/crates/engine-gpu-cuda/src/lib.rs @@ -358,6 +358,10 @@ impl CudaEngine { .with_context(|| "alloc/copy d_index")?; let d_distance = cuda::memory::DeviceBuffer::::zeroed(64) .with_context(|| "alloc d_distance")?; + let d_dbg_y = cuda::memory::DeviceBuffer::::zeroed(64) + .with_context(|| "alloc d_dbg_y")?; + let d_dbg_h = cuda::memory::DeviceBuffer::::zeroed(64) + .with_context(|| "alloc d_dbg_h")?; // Launch G2 kernel // Attempt to populate __constant__ symbols for per-job constants (G2 fast path). @@ -475,6 +479,8 @@ impl CudaEngine { d_found.as_device_ptr(), d_index.as_device_ptr(), d_distance.as_device_ptr(), + d_dbg_y.as_device_ptr(), + d_dbg_h.as_device_ptr(), num_threads as u32, iters_per_thread as u32 )) @@ -490,15 +496,7 @@ impl CudaEngine { .map(|v| v == "1" || v.eq_ignore_ascii_case("true")) .unwrap_or(false) { - if let ( - Ok(mut c_en), - Ok(mut c_idx), - Ok(mut c_dec), - Ok(mut c_y), - Ok(mut c_h), - Ok(mut c_t), - Ok(mut c_th), - ) = ( + if let (Ok(c_en), Ok(c_idx), Ok(c_dec), Ok(c_y), Ok(c_h), Ok(c_t), Ok(c_th)) = ( module.get_global::(CString::new("C_SAMPLER_ENABLE")?.as_c_str()), module.get_global::(CString::new("C_SAMPLER_INDEX")?.as_c_str()), module.get_global::(CString::new("C_SAMPLER_DECISION")?.as_c_str()), @@ -589,10 +587,24 @@ impl CudaEngine { }); } else { // Treat as not found: account coverage and advance like the "not found" path - log::warn!(target: "miner", "CUDA G2: candidate rejected by host re-verification (false positive): idx={k}"); + // Copy debug y/h bytes from device and log first 16 bytes (hex) for diagnosis + let mut dbg_y = [0u8; 64]; + let mut dbg_h = [0u8; 64]; + // Ignore copy errors in logging path; continue accounting regardless + let _ = d_dbg_y.copy_to(&mut dbg_y); + let _ = d_dbg_h.copy_to(&mut dbg_h); + let y_hex = hex::encode(&dbg_y[..16]); + let h_hex = hex::encode(&dbg_h[..16]); + log::warn!( + target: "miner", + "CUDA G2: candidate rejected by host re-verification (false positive): idx={k}, y[0..16]={}, h[0..16]={}", + y_hex, + h_hex + ); #[cfg(feature = "metrics")] { metrics::inc_candidates_false_positive("gpu-cuda"); + metrics::inc_sample_mismatch("gpu-cuda"); } hash_count = hash_count.saturating_add(covered); if covered < total_elems_u64 {