From 751f551141f949340d6b6ab05b6219c0c65532d5 Mon Sep 17 00:00:00 2001 From: shenjack-5600u <3695888@qq.com> Date: Tue, 30 Apr 2024 06:09:08 +0800 Subject: [PATCH] =?UTF-8?q?=E8=A6=81=E9=95=BF=E8=84=91=E5=AD=90=E4=BA=86.p?= =?UTF-8?q?ng?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- miner/src/cacluate.rs | 3 +- miner/src/main.rs | 8 ++--- wgpu-mine/main.wgsl | 18 ++++++++--- wgpu-mine/src/main.rs | 75 +++++++++++++++++++++++++++++++++++-------- 4 files changed, 78 insertions(+), 26 deletions(-) diff --git a/miner/src/cacluate.rs b/miner/src/cacluate.rs index 11be004..13bf491 100644 --- a/miner/src/cacluate.rs +++ b/miner/src/cacluate.rs @@ -109,11 +109,10 @@ pub fn cacl(config: CacluateConfig, id: u64, outfile: &PathBuf) { let xu = crate::evaluate::xuping::XuPing2_0_1015::evaluate(&namer); let xu_qd = crate::evaluate::xuping::XuPing2_0_1015_QD::evaluate(&namer); - if xu < config.qp_expect as f64 && xu_qd < config.qp_expect as f64{ + if xu < config.qp_expect as f64 && xu_qd < config.qp_expect as f64 { continue; } - get_count += 1; info!("Id:{:>15}|{}|{:.4}|{:.4}|{}", i, full_name, xu, xu_qd, namer.get_info()); diff --git a/miner/src/main.rs b/miner/src/main.rs index 7358293..80022c2 100644 --- a/miner/src/main.rs +++ b/miner/src/main.rs @@ -54,11 +54,7 @@ impl Command { qp_expect: self.qp_expect, team: self.team.clone(), report_interval: self.report_interval, - core_affinity: if self.bench { - Some(1 << self.bench_core) - } else { - None - }, + core_affinity: if self.bench { Some(1 << self.bench_core) } else { None }, } } } @@ -84,7 +80,7 @@ pub fn set_thread2core(core: usize) { pub fn set_process_cores(cores: usize) { #[cfg(windows)] unsafe { - use windows_sys::Win32::System::Threading::{SetProcessAffinityMask, GetCurrentProcess}; + use windows_sys::Win32::System::Threading::{GetCurrentProcess, SetProcessAffinityMask}; let process = GetCurrentProcess(); let core_mask = cores; match SetProcessAffinityMask(process, core_mask) { diff --git a/wgpu-mine/main.wgsl b/wgpu-mine/main.wgsl index 085e166..a9918ed 100644 --- a/wgpu-mine/main.wgsl +++ b/wgpu-mine/main.wgsl @@ -1,10 +1,20 @@ @group(0) @binding(0) var first_step: array, 64>; -// 处理第二步的两次 rc4 名称, 输入名字 bytes, 从第一步得到的 val +@group(1) @binding(0) var name_len: array; +@group(1) @binding(1) var name_bytes: array, 256>; + @compute @workgroup_size(16, 16) -fn rc4_name(name_bytes: array, name_len: u32) -> array { - // var val: array = first_step; +// 处理第二步的两次 rc4 名称, 输入名字 bytes, 从第一步得到的 val +// fn rc4_name(@location(0) workname_bytes: array, @location(1) name_len: u32) -> @location(0) array { +// 输入计算着色器的 index, 从全局数据中取对应位置的数据 +fn rc4_name(@builtin(local_invocation_id) compute_pos: vec3) { + // 计算线程的 index + var index: u32 = compute_pos.x + compute_pos.y * 16u; + // 从全局数据中取对应位置的数据 + var name_bytes: array = name_bytes[index]; + var name_len: u32 = name_len[index]; + var val: array = array(); // 把 first_step 的值复制到 val // first_step 内当成连续的内存即可 @@ -34,5 +44,5 @@ fn rc4_name(name_bytes: array, name_len: u32) -> array { } } } - return val; + return; } diff --git a/wgpu-mine/src/main.rs b/wgpu-mine/src/main.rs index fbef3cc..879eec9 100644 --- a/wgpu-mine/src/main.rs +++ b/wgpu-mine/src/main.rs @@ -6,15 +6,41 @@ pub struct Works { /// 队伍名(统一) pub team: String, /// 队伍名(每个任务) - pub names: Vec, + pub names: Vec>, +} + +impl Works { + /// 创建一个新的工作信息 + pub fn new(team: String, names: Vec) -> Self { + // 把输入的名字填充到 256 长度的数组里 + // 如果不够 256 长度, 就用空字符串填充 + let names = { + if names.len() < 256 { + // 先把数组扩容到 256 长度 + let mut names = names; + names.resize(256, "".to_string()); + // 把数组放到一个新的数组里 + vec![names] + } else { + // 如果长度已经是 256 了, 直接截取 + // 不满 256 的部分用空字符串填充 + names + .chunks_exact(256) + .map(|c| { + let mut c = c.to_vec(); + c.resize(256, "".to_string()); + c + }) + .collect::>>() + } + }; + Self { team, names } + } } #[cfg_attr(test, allow(dead_code))] async fn run() { - let works = Works { - team: "team".to_string(), - names: vec!["name1".to_string(), "name2".to_string()], - }; + let works = Works::new("team".to_string(), vec!["name1".to_string(), "name2".to_string()]); let steps = execute_gpu(works).await.unwrap(); @@ -97,10 +123,10 @@ async fn execute_gpu_inner(device: &wgpu::Device, queue: &wgpu::Queue, works: Wo // A storage buffer (can be bound within a bind group and thus available to a shader). // The destination of a copy. // The source of a copy. - let storage_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor { + let uniform_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor { label: Some("team bytes input buffer"), contents: bytemuck::cast_slice(&uniform_val), - usage: wgpu::BufferUsages::STORAGE, + usage: wgpu::BufferUsages::UNIFORM, }); // A bind group defines how buffers are accessed by shaders. @@ -114,21 +140,41 @@ async fn execute_gpu_inner(device: &wgpu::Device, queue: &wgpu::Queue, works: Wo label: None, layout: None, module: &cs_module, - entry_point: "team_bytes", + entry_point: "rc4_name", // constants: &Default::default(), }); // Instantiates the bind group, once again specifying the binding of buffers. - let bind_group_layout = compute_pipeline.get_bind_group_layout(0); - let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor { + let bind_group_layout_0 = compute_pipeline.get_bind_group_layout(0); + let bind_group_0 = device.create_bind_group(&wgpu::BindGroupDescriptor { label: None, - layout: &bind_group_layout, + layout: &bind_group_layout_0, entries: &[wgpu::BindGroupEntry { binding: 0, - resource: storage_buffer.as_entire_binding(), + resource: uniform_buffer.as_entire_binding(), }], }); + let bind_group_layout_1 = compute_pipeline.get_bind_group_layout(1); + let bind_group_1 = device.create_bind_group(&wgpu::BindGroupDescriptor { + label: None, + layout: &bind_group_layout_1, + entries: &[ + wgpu::BindGroupEntry { + binding: 0, + resource: staging_buffer.as_entire_binding(), + }, + // wgpu::BindGroupEntry { + // binding: 1, + // resource: wgpu::BindingResource::Buffer(wgpu::BufferBinding { + // buffer: uniform_buffer, + // offset: 0, + // size, + // }), + // }, + ], + }); + // A command encoder executes one or many pipelines. // It is to WebGPU what a command buffer is to Vulkan. let mut encoder = device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None }); @@ -138,9 +184,10 @@ async fn execute_gpu_inner(device: &wgpu::Device, queue: &wgpu::Queue, works: Wo timestamp_writes: None, }); cpass.set_pipeline(&compute_pipeline); - cpass.set_bind_group(0, &bind_group, &[]); + cpass.set_bind_group(0, &bind_group_0, &[]); + cpass.set_bind_group(1, &bind_group_1, &[]); cpass.insert_debug_marker("compute collatz iterations"); - cpass.dispatch_workgroups(numbers.len() as u32, 1, 1); // Number of cells to run, the (x,y,z) size of item being processed + cpass.dispatch_workgroups(uniform_val.len() as u32, 1, 1); // Number of cells to run, the (x,y,z) size of item being processed } // Sets adds copy operation to command encoder. // Will copy data from storage buffer on GPU to staging buffer on CPU.