gpu runs but doesn't mine
This commit is contained in:
illuzen
3
Cargo.lock
generated
3
Cargo.lock
generated
@@ -560,9 +560,11 @@ dependencies = [
|
||||
"anyhow",
|
||||
"bytemuck",
|
||||
"criterion",
|
||||
"engine-cpu",
|
||||
"futures",
|
||||
"hex",
|
||||
"log",
|
||||
"pow-core",
|
||||
"primitive-types 0.13.1",
|
||||
"qp-plonky2 1.1.1",
|
||||
"qp-plonky2-field 1.1.1",
|
||||
@@ -1428,6 +1430,7 @@ dependencies = [
|
||||
"anyhow",
|
||||
"crossbeam-channel",
|
||||
"engine-cpu",
|
||||
"engine-gpu",
|
||||
"hex",
|
||||
"log",
|
||||
"metrics",
|
||||
|
||||
@@ -18,6 +18,8 @@ baseline = []
|
||||
simd-poseidon2 = []
|
||||
|
||||
[dependencies]
|
||||
engine-cpu = { path = "../engine-cpu" }
|
||||
pow-core = { path = "../pow-core" }
|
||||
primitive-types = { workspace = true }
|
||||
log = { workspace = true }
|
||||
thiserror = { workspace = true }
|
||||
|
||||
271
crates/engine-gpu/src/lib.rs
Normal file
271
crates/engine-gpu/src/lib.rs
Normal file
@@ -0,0 +1,271 @@
|
||||
#![deny(rust_2018_idioms)]
|
||||
#![forbid(unsafe_code)]
|
||||
|
||||
use engine_cpu::{Candidate, EngineStatus, FoundOrigin, MinerEngine, Range};
|
||||
use futures::executor::block_on;
|
||||
use pow_core::JobContext;
|
||||
use primitive_types::U512;
|
||||
use std::sync::atomic::{AtomicBool, Ordering};
|
||||
use wgpu::util::DeviceExt;
|
||||
|
||||
pub struct GpuEngine {
|
||||
device: wgpu::Device,
|
||||
queue: wgpu::Queue,
|
||||
pipeline: wgpu::ComputePipeline,
|
||||
bind_group_layout: wgpu::BindGroupLayout,
|
||||
}
|
||||
|
||||
impl GpuEngine {
|
||||
pub fn new() -> Self {
|
||||
block_on(Self::init()).expect("Failed to initialize GPU engine")
|
||||
}
|
||||
|
||||
async fn init() -> Result<Self, Box<dyn std::error::Error>> {
|
||||
let instance = wgpu::Instance::new(&wgpu::InstanceDescriptor {
|
||||
backends: wgpu::Backends::PRIMARY,
|
||||
..Default::default()
|
||||
});
|
||||
|
||||
let adapter = instance
|
||||
.request_adapter(&wgpu::RequestAdapterOptions {
|
||||
power_preference: wgpu::PowerPreference::HighPerformance,
|
||||
..Default::default()
|
||||
})
|
||||
.await
|
||||
.map_err(|e| format!("No suitable GPU adapter found: {:?}", e))?;
|
||||
|
||||
let (device, queue) = adapter
|
||||
.request_device(&wgpu::DeviceDescriptor {
|
||||
label: Some("Mining Device"),
|
||||
required_features: wgpu::Features::empty(),
|
||||
required_limits: wgpu::Limits::default(),
|
||||
memory_hints: Default::default(),
|
||||
..Default::default()
|
||||
})
|
||||
.await?;
|
||||
|
||||
let shader_source = include_str!("mining.wgsl");
|
||||
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
|
||||
label: Some("Mining Shader"),
|
||||
source: wgpu::ShaderSource::Wgsl(shader_source.into()),
|
||||
});
|
||||
|
||||
let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
|
||||
label: Some("Mining Pipeline"),
|
||||
layout: None,
|
||||
module: &shader,
|
||||
entry_point: Some("mining_main"),
|
||||
compilation_options: Default::default(),
|
||||
cache: None,
|
||||
});
|
||||
|
||||
let bind_group_layout = pipeline.get_bind_group_layout(0);
|
||||
|
||||
Ok(Self {
|
||||
device,
|
||||
queue,
|
||||
pipeline,
|
||||
bind_group_layout,
|
||||
})
|
||||
}
|
||||
|
||||
fn run_batch(&self, ctx: &JobContext, start_nonce: U512, batch_size: u32) -> Option<Candidate> {
|
||||
// Prepare buffers
|
||||
let mut header_u32s = [0u32; 8];
|
||||
for i in 0..8 {
|
||||
let chunk = &ctx.header[i * 4..(i + 1) * 4];
|
||||
header_u32s[i] = u32::from_le_bytes([chunk[0], chunk[1], chunk[2], chunk[3]]);
|
||||
}
|
||||
|
||||
let start_nonce_bytes = start_nonce.to_little_endian();
|
||||
let mut start_nonce_u32s = [0u32; 16];
|
||||
for i in 0..16 {
|
||||
let chunk = &start_nonce_bytes[i * 4..(i + 1) * 4];
|
||||
start_nonce_u32s[i] = u32::from_le_bytes([chunk[0], chunk[1], chunk[2], chunk[3]]);
|
||||
}
|
||||
|
||||
let target_bytes = ctx.difficulty.to_little_endian();
|
||||
let mut target_u32s = [0u32; 16];
|
||||
for i in 0..16 {
|
||||
let chunk = &target_bytes[i * 4..(i + 1) * 4];
|
||||
target_u32s[i] = u32::from_le_bytes([chunk[0], chunk[1], chunk[2], chunk[3]]);
|
||||
}
|
||||
|
||||
let header_buffer = self
|
||||
.device
|
||||
.create_buffer_init(&wgpu::util::BufferInitDescriptor {
|
||||
label: Some("Header Buffer"),
|
||||
contents: bytemuck::cast_slice(&header_u32s),
|
||||
usage: wgpu::BufferUsages::STORAGE,
|
||||
});
|
||||
|
||||
let start_nonce_buffer =
|
||||
self.device
|
||||
.create_buffer_init(&wgpu::util::BufferInitDescriptor {
|
||||
label: Some("Start Nonce Buffer"),
|
||||
contents: bytemuck::cast_slice(&start_nonce_u32s),
|
||||
usage: wgpu::BufferUsages::STORAGE,
|
||||
});
|
||||
|
||||
let target_buffer = self
|
||||
.device
|
||||
.create_buffer_init(&wgpu::util::BufferInitDescriptor {
|
||||
label: Some("Target Buffer"),
|
||||
contents: bytemuck::cast_slice(&target_u32s),
|
||||
usage: wgpu::BufferUsages::STORAGE,
|
||||
});
|
||||
|
||||
// Results: [flag (1), nonce (16), hash (16)] = 33 u32s
|
||||
let results_size = (1 + 16 + 16) * 4;
|
||||
let results_buffer = self.device.create_buffer(&wgpu::BufferDescriptor {
|
||||
label: Some("Results Buffer"),
|
||||
size: results_size,
|
||||
usage: wgpu::BufferUsages::STORAGE
|
||||
| wgpu::BufferUsages::COPY_SRC
|
||||
| wgpu::BufferUsages::COPY_DST,
|
||||
mapped_at_creation: false,
|
||||
});
|
||||
|
||||
// Initialize results to 0
|
||||
let zeros = vec![0u8; results_size as usize];
|
||||
self.queue.write_buffer(&results_buffer, 0, &zeros);
|
||||
|
||||
let bind_group = self.device.create_bind_group(&wgpu::BindGroupDescriptor {
|
||||
label: Some("Mining Bind Group"),
|
||||
layout: &self.bind_group_layout,
|
||||
entries: &[
|
||||
wgpu::BindGroupEntry {
|
||||
binding: 0,
|
||||
resource: results_buffer.as_entire_binding(),
|
||||
},
|
||||
wgpu::BindGroupEntry {
|
||||
binding: 1,
|
||||
resource: header_buffer.as_entire_binding(),
|
||||
},
|
||||
wgpu::BindGroupEntry {
|
||||
binding: 2,
|
||||
resource: start_nonce_buffer.as_entire_binding(),
|
||||
},
|
||||
wgpu::BindGroupEntry {
|
||||
binding: 3,
|
||||
resource: target_buffer.as_entire_binding(),
|
||||
},
|
||||
],
|
||||
});
|
||||
|
||||
let mut encoder = self
|
||||
.device
|
||||
.create_command_encoder(&wgpu::CommandEncoderDescriptor {
|
||||
label: Some("Mining Encoder"),
|
||||
});
|
||||
|
||||
{
|
||||
let mut cpass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
|
||||
label: Some("Mining Compute Pass"),
|
||||
timestamp_writes: None,
|
||||
});
|
||||
cpass.set_pipeline(&self.pipeline);
|
||||
cpass.set_bind_group(0, &bind_group, &[]);
|
||||
let workgroups = (batch_size + 255) / 256;
|
||||
cpass.dispatch_workgroups(workgroups, 1, 1);
|
||||
}
|
||||
|
||||
let staging_buffer = self.device.create_buffer(&wgpu::BufferDescriptor {
|
||||
label: Some("Staging Buffer"),
|
||||
size: results_size,
|
||||
usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST,
|
||||
mapped_at_creation: false,
|
||||
});
|
||||
|
||||
encoder.copy_buffer_to_buffer(&results_buffer, 0, &staging_buffer, 0, results_size);
|
||||
|
||||
self.queue.submit(Some(encoder.finish()));
|
||||
|
||||
let buffer_slice = staging_buffer.slice(..);
|
||||
let (sender, receiver) = futures::channel::oneshot::channel();
|
||||
buffer_slice.map_async(wgpu::MapMode::Read, move |v| sender.send(v).unwrap());
|
||||
|
||||
let _ = self.device.poll(wgpu::PollType::Wait {
|
||||
submission_index: None,
|
||||
timeout: None,
|
||||
});
|
||||
|
||||
if let Ok(Ok(())) = block_on(receiver) {
|
||||
let data = buffer_slice.get_mapped_range();
|
||||
let result_u32s: &[u32] = bytemuck::cast_slice(&data);
|
||||
|
||||
if result_u32s[0] != 0 {
|
||||
let mut nonce_bytes = [0u8; 64];
|
||||
for i in 0..16 {
|
||||
let bytes = result_u32s[1 + i].to_le_bytes();
|
||||
nonce_bytes[i * 4..(i + 1) * 4].copy_from_slice(&bytes);
|
||||
}
|
||||
let nonce = U512::from_little_endian(&nonce_bytes);
|
||||
|
||||
let mut hash_bytes = [0u8; 64];
|
||||
for i in 0..16 {
|
||||
let bytes = result_u32s[17 + i].to_le_bytes();
|
||||
hash_bytes[i * 4..(i + 1) * 4].copy_from_slice(&bytes);
|
||||
}
|
||||
let hash = U512::from_little_endian(&hash_bytes);
|
||||
|
||||
let work = nonce.to_big_endian();
|
||||
|
||||
return Some(Candidate { nonce, work, hash });
|
||||
}
|
||||
}
|
||||
|
||||
None
|
||||
}
|
||||
}
|
||||
|
||||
impl MinerEngine for GpuEngine {
|
||||
fn name(&self) -> &'static str {
|
||||
"gpu-wgpu"
|
||||
}
|
||||
|
||||
fn prepare_context(&self, header_hash: [u8; 32], difficulty: U512) -> JobContext {
|
||||
JobContext::new(header_hash, difficulty)
|
||||
}
|
||||
|
||||
fn search_range(&self, ctx: &JobContext, range: Range, cancel: &AtomicBool) -> EngineStatus {
|
||||
let mut current_start = range.start;
|
||||
let mut hash_count = 0;
|
||||
let batch_size = 65536; // 256 * 256
|
||||
|
||||
if range.start > range.end {
|
||||
return EngineStatus::Exhausted { hash_count: 0 };
|
||||
}
|
||||
|
||||
while current_start <= range.end {
|
||||
if cancel.load(Ordering::Relaxed) {
|
||||
return EngineStatus::Cancelled { hash_count };
|
||||
}
|
||||
|
||||
let remaining = range.end - current_start + 1;
|
||||
let current_batch_size = if remaining < U512::from(batch_size) {
|
||||
remaining.as_u64() as u32
|
||||
} else {
|
||||
batch_size
|
||||
};
|
||||
|
||||
if let Some(candidate) = self.run_batch(ctx, current_start, current_batch_size) {
|
||||
// We found a candidate!
|
||||
// The hash_count is approximate since we don't know exactly which thread found it
|
||||
// without reading back the index, but we can just say we did the whole batch.
|
||||
// Or we can calculate it from candidate.nonce - current_start.
|
||||
let found_offset = (candidate.nonce - current_start).as_u64();
|
||||
return EngineStatus::Found {
|
||||
candidate,
|
||||
hash_count: hash_count + found_offset + 1,
|
||||
origin: FoundOrigin::GpuG1,
|
||||
};
|
||||
}
|
||||
|
||||
hash_count += current_batch_size as u64;
|
||||
current_start = current_start + U512::from(current_batch_size);
|
||||
}
|
||||
|
||||
EngineStatus::Exhausted { hash_count }
|
||||
}
|
||||
}
|
||||
@@ -190,7 +190,7 @@ const MDS_MATRIX_DIAG_12: array<array<u32, 2>, 12> = array<array<u32, 2>, 12>(
|
||||
);
|
||||
|
||||
// Storage buffers
|
||||
@group(0) @binding(0) var<storage, read_write> results: array<u32>;
|
||||
@group(0) @binding(0) var<storage, read_write> results: array<atomic<u32>>;
|
||||
@group(0) @binding(1) var<storage, read> header: array<u32, 8>; // 32 bytes
|
||||
@group(0) @binding(2) var<storage, read> start_nonce: array<u32, 16>; // 64 bytes
|
||||
@group(0) @binding(3) var<storage, read> difficulty_target: array<u32, 16>; // 64 bytes (U512 target)
|
||||
@@ -879,4 +879,65 @@ fn is_below_target(hash: array<u32, 16>, difficulty_tgt: array<u32, 16>) -> bool
|
||||
return false; // Equal, not below
|
||||
}
|
||||
|
||||
// No main function - individual tests will create their own entry points
|
||||
// Main mining kernel
|
||||
@compute @workgroup_size(256)
|
||||
fn mining_main(@builtin(global_invocation_id) global_id: vec3<u32>) {
|
||||
// If solution already found, exit early
|
||||
if (atomicLoad(&results[0]) != 0u) {
|
||||
return;
|
||||
}
|
||||
|
||||
let index = global_id.x;
|
||||
|
||||
// Compute current nonce = start_nonce + index
|
||||
var current_nonce: array<u32, 16>;
|
||||
var carry: u32 = 0u;
|
||||
|
||||
// First limb addition with index
|
||||
let val0 = start_nonce[0];
|
||||
let sum0 = val0 + index;
|
||||
current_nonce[0] = sum0;
|
||||
carry = select(0u, 1u, sum0 < val0);
|
||||
|
||||
// Subsequent limbs with carry propagation
|
||||
for (var i = 1u; i < 16u; i++) {
|
||||
let val = start_nonce[i];
|
||||
let sum = val + carry;
|
||||
current_nonce[i] = sum;
|
||||
carry = select(0u, 1u, sum < val);
|
||||
}
|
||||
|
||||
// Construct input (96 bytes = 24 u32s)
|
||||
// Header (32 bytes = 8 u32s) followed by Nonce (64 bytes = 16 u32s)
|
||||
var input: array<u32, 24>;
|
||||
for (var i = 0u; i < 8u; i++) {
|
||||
input[i] = header[i];
|
||||
}
|
||||
// Nonce needs to be Big Endian in the byte stream for hashing.
|
||||
// current_nonce is Little Endian words.
|
||||
for (var i = 0u; i < 16u; i++) {
|
||||
let val = current_nonce[15u - i];
|
||||
// Reverse bytes
|
||||
let rev = ((val & 0xFFu) << 24u) |
|
||||
((val & 0xFF00u) << 8u) |
|
||||
((val & 0xFF0000u) >> 8u) |
|
||||
((val & 0xFF000000u) >> 24u);
|
||||
input[8u + i] = rev;
|
||||
}
|
||||
|
||||
// Hash
|
||||
let hash = double_hash(input);
|
||||
|
||||
// Check target
|
||||
if (is_below_target(hash, difficulty_target)) {
|
||||
// Try to claim the solution
|
||||
if (atomicExchange(&results[0], 1u) == 0u) {
|
||||
// We won! Write nonce and hash
|
||||
// results layout: [0]=found, [1..16]=nonce, [17..32]=hash
|
||||
for (var i = 0u; i < 16u; i++) {
|
||||
atomicStore(&results[1u + i], current_nonce[i]);
|
||||
atomicStore(&results[17u + i], hash[i]);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@@ -7,6 +7,7 @@ description = "CLI binary to run the Quantus External Miner service"
|
||||
|
||||
[features]
|
||||
default = []
|
||||
gpu = ["miner-service/gpu"]
|
||||
|
||||
[dependencies]
|
||||
miner-service = { path = "../miner-service" }
|
||||
|
||||
@@ -40,8 +40,7 @@ struct Args {
|
||||
manip_throttle_cap: Option<u64>,
|
||||
|
||||
/// Mining engine to use (default: cpu-fast).
|
||||
/// Options: cpu-baseline, cpu-fast, cpu-chain-manipulator, gpu-cuda, gpu-opencl
|
||||
/// Note: GPU engines are currently unimplemented and will return a clear error at runtime.
|
||||
/// Options: cpu-baseline, cpu-fast, cpu-chain-manipulator, gpu
|
||||
#[arg(long, env = "MINER_ENGINE", value_enum, default_value_t = EngineCli::CpuFast)]
|
||||
engine: EngineCli,
|
||||
|
||||
@@ -98,10 +97,8 @@ enum EngineCli {
|
||||
CpuFast,
|
||||
/// Throttling CPU engine that slows per block to help reduce difficulty
|
||||
CpuChainManipulator,
|
||||
/// CUDA GPU engine (unimplemented; selecting will return an error)
|
||||
GpuCuda,
|
||||
/// OpenCL GPU engine (unimplemented; selecting will return an error)
|
||||
GpuOpencl,
|
||||
/// GPU engine (WGPU based)
|
||||
Gpu,
|
||||
}
|
||||
|
||||
impl From<EngineCli> for EngineSelection {
|
||||
@@ -110,8 +107,7 @@ impl From<EngineCli> for EngineSelection {
|
||||
EngineCli::CpuBaseline => EngineSelection::CpuBaseline,
|
||||
EngineCli::CpuFast => EngineSelection::CpuFast,
|
||||
EngineCli::CpuChainManipulator => EngineSelection::CpuChainManipulator,
|
||||
EngineCli::GpuCuda => EngineSelection::GpuCuda,
|
||||
EngineCli::GpuOpencl => EngineSelection::GpuOpenCl,
|
||||
EngineCli::Gpu => EngineSelection::Gpu,
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@@ -9,6 +9,8 @@ description = "Service layer: HTTP API, job orchestration, and engine abstractio
|
||||
default = ["cpu", "metrics"]
|
||||
# Enable CPU engine by default.
|
||||
cpu = ["engine-cpu"]
|
||||
# Enable GPU engine.
|
||||
gpu = ["engine-gpu"]
|
||||
# Optional metrics/observability (Prometheus endpoint).
|
||||
metrics = [
|
||||
"dep:metrics",
|
||||
@@ -35,5 +37,6 @@ quantus-miner-api = { workspace = true }
|
||||
# Local crates
|
||||
pow-core = { path = "../pow-core" }
|
||||
engine-cpu = { path = "../engine-cpu", optional = true }
|
||||
engine-gpu = { path = "../engine-gpu", optional = true }
|
||||
metrics = { path = "../metrics", optional = true }
|
||||
miner-telemetry = { path = "../miner-telemetry" }
|
||||
|
||||
@@ -45,8 +45,7 @@ pub enum EngineSelection {
|
||||
CpuBaseline,
|
||||
CpuFast,
|
||||
CpuChainManipulator,
|
||||
GpuCuda,
|
||||
GpuOpenCl,
|
||||
Gpu,
|
||||
}
|
||||
|
||||
impl Default for ServiceConfig {
|
||||
@@ -71,8 +70,7 @@ impl fmt::Display for ServiceConfig {
|
||||
EngineSelection::CpuBaseline => "cpu-baseline",
|
||||
EngineSelection::CpuFast => "cpu-fast",
|
||||
EngineSelection::CpuChainManipulator => "cpu-chain-manipulator",
|
||||
EngineSelection::GpuCuda => "gpu-cuda",
|
||||
EngineSelection::GpuOpenCl => "gpu-opencl",
|
||||
EngineSelection::Gpu => "gpu",
|
||||
};
|
||||
write!(
|
||||
f,
|
||||
@@ -1194,32 +1192,19 @@ pub async fn run(config: ServiceConfig) -> anyhow::Result<()> {
|
||||
}
|
||||
Arc::new(eng)
|
||||
}
|
||||
EngineSelection::GpuCuda => {
|
||||
#[cfg(feature = "cuda")]
|
||||
EngineSelection::Gpu => {
|
||||
#[cfg(feature = "gpu")]
|
||||
{
|
||||
let eng = engine_gpu_cuda::CudaEngine::new();
|
||||
if !eng.cuda_available() {
|
||||
log::error!("Requested engine gpu-cuda but CUDA driver/device is not available or initialization failed. Build has 'cuda' feature but runtime support is missing.");
|
||||
return Err(anyhow::anyhow!(
|
||||
"engine 'gpu-cuda' unavailable at runtime (CUDA init failed)"
|
||||
));
|
||||
}
|
||||
Arc::new(eng)
|
||||
Arc::new(engine_gpu::GpuEngine::new())
|
||||
}
|
||||
#[cfg(not(feature = "cuda"))]
|
||||
#[cfg(not(feature = "gpu"))]
|
||||
{
|
||||
log::error!("Requested engine gpu-cuda, but this binary was built without the 'cuda' feature. Rebuild miner-service with --features cuda.");
|
||||
log::error!("Requested engine gpu, but this binary was built without the 'gpu' feature. Rebuild miner-service with --features gpu.");
|
||||
return Err(anyhow::anyhow!(
|
||||
"engine 'gpu-cuda' not built (missing 'cuda' feature)"
|
||||
"engine 'gpu' not built (missing 'gpu' feature)"
|
||||
));
|
||||
}
|
||||
}
|
||||
EngineSelection::GpuOpenCl => {
|
||||
log::error!("Requested engine gpu-opencl is not implemented yet. Use a CPU engine (cpu-fast or cpu-baseline) for now.");
|
||||
return Err(anyhow::anyhow!(
|
||||
"engine 'gpu-opencl' is not implemented yet"
|
||||
));
|
||||
}
|
||||
};
|
||||
log::info!("Using engine: {}", engine.name());
|
||||
log::info!("Service configuration: {config}");
|
||||
|
||||
Reference in New Issue
Block a user