diff --git a/crates/engine-gpu-cuda/src/kernels/qpow_kernel.cu b/crates/engine-gpu-cuda/src/kernels/qpow_kernel.cu index 4db62ca..ab94228 100644 --- a/crates/engine-gpu-cuda/src/kernels/qpow_kernel.cu +++ b/crates/engine-gpu-cuda/src/kernels/qpow_kernel.cu @@ -323,9 +323,18 @@ __device__ __constant__ uint64_t C_R2[8]; __device__ __constant__ uint64_t C_MHAT[8]; __device__ __constant__ uint64_t C_N0_INV; __device__ __constant__ int C_CONSTS_READY; -// Optional device-side target/threshold in constant memory (set by host when available) +// Optional constant-memory target/threshold for device compare __device__ __constant__ uint64_t C_TARGET[8]; __device__ __constant__ uint64_t C_THRESH[8]; +// Optional sampler controls/output (host may read these symbols when enabled) +__device__ __constant__ int C_SAMPLER_ENABLE; +__device__ uint8_t C_SAMPLER_Y_BE[64]; +__device__ uint8_t C_SAMPLER_H_BE[64]; +__device__ uint8_t C_SAMPLER_TARGET_BE[64]; +__device__ uint8_t C_SAMPLER_THRESH_BE[64]; +__device__ uint32_t C_SAMPLER_INDEX; +__device__ uint32_t C_SAMPLER_DECISION; +/* removed duplicate C_TARGET/C_THRESH definitions */ // Load/store helpers (little-endian) __device__ __forceinline__ uint64_t load64_le(const uint8_t* p) { @@ -348,6 +357,16 @@ __device__ __forceinline__ void store64_le(uint8_t* p, uint64_t v) { p[6] = (uint8_t)(v >> 48); p[7] = (uint8_t)(v >> 56); } +__device__ __forceinline__ void store64_be(uint8_t* p, uint64_t v) { + p[0] = (uint8_t)(v >> 56); + p[1] = (uint8_t)(v >> 48); + p[2] = (uint8_t)(v >> 40); + p[3] = (uint8_t)(v >> 32); + p[4] = (uint8_t)(v >> 24); + p[5] = (uint8_t)(v >> 16); + p[6] = (uint8_t)(v >> 8); + p[7] = (uint8_t)(v); +} // Keccak-f[1600] permutation __device__ __forceinline__ void keccak_f1600(uint64_t s[25]) { @@ -430,18 +449,10 @@ __device__ __forceinline__ void sha3_512_64bytes(const uint8_t in_be64[64], uint // Permute keccak_f1600(s); - // Squeeze 64 bytes (8 lanes) - uint8_t out_le[64]; -#pragma unroll + // Squeeze 64 bytes (8 lanes) into big-endian digest bytes (match host SHA3 output) + #pragma unroll for (int i = 0; i < 8; ++i) { - store64_le(&out_le[i * 8], s[i]); - } - - // The standard digest byte order expected by host libraries is the emitted byte stream. - // We keep it as-is; host target/threshold bytes should be in the same orientation to XOR/compare. -#pragma unroll - for (int i = 0; i < 64; ++i) { - out_be64[i] = out_le[i]; + store64_be(&out_be64[i * 8], s[i]); } } diff --git a/crates/engine-gpu-cuda/src/lib.rs b/crates/engine-gpu-cuda/src/lib.rs index b4ff12e..d7bade0 100644 --- a/crates/engine-gpu-cuda/src/lib.rs +++ b/crates/engine-gpu-cuda/src/lib.rs @@ -399,6 +399,53 @@ impl CudaEngine { c_n0.copy_from(&(n0_inv as u64))?; c_target.copy_from(&target_limbs)?; c_thresh.copy_from(&thresh_limbs)?; + // Optional sampler: initialize when env MINER_CUDA_SAMPLER=1 + if std::env::var("MINER_CUDA_SAMPLER") + .map(|v| v == "1" || v.eq_ignore_ascii_case("true")) + .unwrap_or(false) + { + if let ( + Ok(mut c_samp_en), + Ok(mut c_samp_y), + Ok(mut c_samp_h), + Ok(mut c_samp_target_be), + Ok(mut c_samp_thresh_be), + Ok(mut c_samp_index), + Ok(mut c_samp_dec), + ) = ( + module.get_global::( + CString::new("C_SAMPLER_ENABLE")?.as_c_str(), + ), + module.get_global::<[u8; 64]>( + CString::new("C_SAMPLER_Y_BE")?.as_c_str(), + ), + module.get_global::<[u8; 64]>( + CString::new("C_SAMPLER_H_BE")?.as_c_str(), + ), + module.get_global::<[u8; 64]>( + CString::new("C_SAMPLER_TARGET_BE")?.as_c_str(), + ), + module.get_global::<[u8; 64]>( + CString::new("C_SAMPLER_THRESH_BE")?.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(), + ), + ) { + let one_i32: i32 = 1; + c_samp_en.copy_from(&one_i32)?; + // Initialize sampler buffers (copy target/threshold; zero y/h/index/decision) + c_samp_target_be.copy_from(&target_be)?; + c_samp_thresh_be.copy_from(&threshold_be)?; + c_samp_index.copy_from(&0u32)?; + c_samp_dec.copy_from(&0u32)?; + let zero64 = [0u8; 64]; + c_samp_y.copy_from(&zero64)?; + c_samp_h.copy_from(&zero64)?; + } + } // Flag ready = 1 let one: i32 = 1; c_ready.copy_from(&one)?; @@ -485,11 +532,6 @@ impl CudaEngine { metrics::inc_candidates_false_positive("gpu-cuda"); } hash_count = hash_count.saturating_add(covered); - #[cfg(feature = "metrics")] - { - metrics::inc_candidates_false_positive("gpu-cuda"); - } - hash_count = hash_count.saturating_add(covered); if covered < total_elems_u64 { return Ok(EngineStatus::Exhausted { hash_count }); } diff --git a/crates/metrics/src/lib.rs b/crates/metrics/src/lib.rs index 29b1255..72cd3d2 100644 --- a/crates/metrics/src/lib.rs +++ b/crates/metrics/src/lib.rs @@ -230,6 +230,21 @@ static CANDIDATES_FALSE_POSITIVE_TOTAL: Lazy = Lazy::new(|| { c }); +static SAMPLE_MISMATCH_TOTAL: Lazy = Lazy::new(|| { + let c = IntCounterVec::new( + opts!( + "miner_sample_mismatch_total", + "Total decision parity mismatches between engine and host per engine" + ), + &["engine"], + ) + .expect("create miner_sample_mismatch_total"); + REGISTRY + .register(Box::new(c.clone())) + .expect("register miner_sample_mismatch_total"); + c +}); + static THREAD_HASHES_TOTAL: Lazy = Lazy::new(|| { let c = IntCounterVec::new( opts!( @@ -344,6 +359,10 @@ pub fn inc_candidates_false_positive(engine: &str) { .inc(); } +pub fn inc_sample_mismatch(engine: &str) { + SAMPLE_MISMATCH_TOTAL.with_label_values(&[engine]).inc(); +} + pub fn inc_thread_hashes(engine: &str, job_id: &str, thread_id: &str, n: u64) { THREAD_HASHES_TOTAL .with_label_values(&[engine, job_id, thread_id])