blackbeard/miner
1
0
Fork 0
Code Issues 2 Pull Requests Actions Packages Projects Releases 2 Wiki Activity

Add optional CUDA sampler support and big-endian SHA3 output

Browse Source 
- Introduce constant memory variables and device buffers for sampler -
Initialize sampler in CudaEngine when MINER_CUDA_SAMPLER=1 - Change SHA3
kernel to output digest bytes in big-endian format - Add sample mismatch
metric counter
This commit is contained in:
rob thijssen
2025-09-13 04:02:06 +03:00
parent 191cc799b7
commit 0aaacdf85d
3 changed files with 89 additions and 17 deletions
Download Patch File Download Diff File Expand all files Collapse all files
crates/engine-gpu-cuda/src/kernels/qpow_kernel.cu
+23 -12
View File
@@ -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]);
}
}
crates/engine-gpu-cuda/src/lib.rs
+47 -5
View File
@@ -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::<i32>(
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::<u32>(CString::new("C_SAMPLER_INDEX")?.as_c_str()),
module.get_global::<u32>(
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 });
}
crates/metrics/src/lib.rs
+19
View File
@@ -230,6 +230,21 @@ static CANDIDATES_FALSE_POSITIVE_TOTAL: Lazy<IntCounterVec> = Lazy::new(|| {
c
});
static SAMPLE_MISMATCH_TOTAL: Lazy<IntCounterVec> = 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<IntCounterVec> = 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])