diff options
Diffstat (limited to 'src')
| -rw-r--r-- | src/main.rs | 489 | ||||
| -rw-r--r-- | src/shader.wgsl | 208 | ||||
| -rw-r--r-- | src/shader_own.wgsl | 226 |
3 files changed, 923 insertions, 0 deletions
diff --git a/src/main.rs b/src/main.rs new file mode 100644 index 0000000..54d9610 --- /dev/null +++ b/src/main.rs @@ -0,0 +1,489 @@ +use std::env::{self, Args}; +use std::mem; +use std::num::NonZeroU64; +use std::ops::RangeBounds; +use std::time::Instant; + +use wgpu::{BufferUsages, SubmissionIndex, include_wgsl}; +use wgpu::util::DeviceExt; +use wgpu::wgc::command; + + + +const PALETTE: &[u8] = include_bytes!("../palette.txt"); +const PALETTE_SIZE: usize = PALETTE.len(); + +const BATCH_SIZE : u64 = (1 << 16) - 64; +const MAX_ITER : u64 = 1 << 32; +// const BATCH_SIZE : u64 = 8; +// const MAX_ITER : u64 = 2; + +const INPUT_SIZE:u64 = 512 / 8 * 8; +const INSIZE_SIZE:u64= 32 / 8; +const OUTPUT_SIZE:u64= 256 / 8; +const INPUT_BUF_SIZE :u64= INPUT_SIZE * BATCH_SIZE; +const SIZES_BUF_SIZE :u64= INSIZE_SIZE * BATCH_SIZE; +const OUTPUT_BUF_SIZE:u64= OUTPUT_SIZE * BATCH_SIZE; + +fn get_arg(args: &mut Args, pname: &String , name: &'static str) -> String { + args.next().unwrap_or_else(|| panic!("Argument missing: {}\nUsage: {} value prev bits", name, pname)).clone() +} + + +struct NonceIter { + digits: Vec<usize>, +} + +impl NonceIter { + fn new_with_capacity(capacity: usize) -> Self { + Self { + digits: Vec::with_capacity(capacity) + } + } +} + +impl Iterator for NonceIter { + type Item = NonceElement; + fn next(&mut self) -> Option<Self::Item> { + let mut carry = 1; + let mut cursor = 0; + while carry > 0 { + if let Some(d) = self.digits.get_mut(cursor) { + if *d + carry >= PALETTE_SIZE { + *d = (*d + carry) % PALETTE_SIZE; + } else { + *d += carry; + carry = 0; + } + } else { + self.digits.push(carry); + carry = 0; + } + cursor += 1; + } + Some(NonceElement::from_digits(&self.digits)) + } + +} + +#[derive(Debug)] +struct NonceElement { + pub bytes: Vec<u8> +} + +impl NonceElement { + fn from_digits(digits: &Vec<usize>) -> Self { + let mut bytes = Vec::with_capacity(digits.len()); + for d in digits { + bytes.push(PALETTE[*d]) + } + Self { + bytes + } + } + fn to_hash_input(&self, value: &String, prev: &String) -> String { + let mut out = String::with_capacity(self.bytes.len() + value.len() + prev.len()); + out += value; + out += prev; + out += str::from_utf8(self.bytes.as_slice()).expect("Invalid bytes in pallette"); + out + } +} + + +/// Swap chain. We Have Swap Chains At Home Edition. +/// When given an index, returns the first tuple entry if the index is even, and the second if it's +/// odd. +fn swaptuple_get<T>(tup :&(T,T), i: u64) -> &T { + if i.is_multiple_of(2){ &tup.0 } else { &tup.1 } +} +/// Swap chain. We Have Swap Chains At Home Edition. Mutable Edition. +/// When given an index, returns the first tuple entry if the index is even, and the second if it's +/// odd, but now mutable. +fn swaptuple_get_mut<T>(tup :& mut (T,T), i: u64) -> & mut T { + if i.is_multiple_of(2) { &mut tup.0 } else { &mut tup.1 } +} + +fn main() { + let mut args = env::args(); + + let pname: String = args.next().expect("Program name should always be included."); + let value: String = get_arg(&mut args, &pname, "value"); + let prev: String = get_arg(&mut args, &pname, "prev"); + let bits: u16 = get_arg(&mut args, &pname, "bits").parse().unwrap(); + + // We first initialize an wgpu `Instance`, which contains any "global" state wgpu needs. + // + // This is what loads the vulkan/dx12/metal/opengl libraries. + env_logger::init(); + + let instance = wgpu::Instance::new(&wgpu::InstanceDescriptor::from_env_or_default()); + let areq = wgpu::RequestAdapterOptions {power_preference: wgpu::PowerPreference::HighPerformance, ..Default::default()}; + let adapter = + pollster::block_on(instance.request_adapter(&areq)) + .expect("Failed to create adapter"); + + // Print out some basic information about the adapter. + println!("Running on Adapter: {:#?}", adapter.get_info()); + + // Check to see if the adapter supports compute shaders. While WebGPU guarantees support for + // compute shaders, wgpu supports a wider range of devices through the use of "downlevel" devices. + let downlevel_capabilities = adapter.get_downlevel_capabilities(); + if !downlevel_capabilities + .flags + .contains(wgpu::DownlevelFlags::COMPUTE_SHADERS) + { + panic!("Adapter does not support compute shaders"); + } + + // We then create a `Device` and a `Queue` from the `Adapter`. + // + // The `Device` is used to create and manage GPU resources. + // The `Queue` is a queue used to submit work for the GPU to process. + let (device, queue) = pollster::block_on(adapter.request_device(&wgpu::DeviceDescriptor { + label: None, + required_features: wgpu::Features::empty(), + required_limits: wgpu::Limits::downlevel_defaults(), + experimental_features: wgpu::ExperimentalFeatures::disabled(), + memory_hints: wgpu::MemoryHints::MemoryUsage, + trace: wgpu::Trace::Off, + })) + .expect("Failed to create device"); + + // Create a shader module from our shader code. This will parse and validate the shader. + // + // `include_wgsl` is a macro provided by wgpu like `include_str` which constructs a ShaderModuleDescriptor. + // If you want to load shaders differently, you can construct the ShaderModuleDescriptor manually. + let module = device.create_shader_module(wgpu::include_wgsl!("shader_own.wgsl")); + + let mut iter = NonceIter::new_with_capacity((512 - prev.len() - value.len())); + + // Create a buffer with the data we want to process on the GPU. + // + // `create_buffer_init` is a utility provided by `wgpu::util::DeviceExt` which simplifies creating + // a buffer with some initial data. + // + // We use the `bytemuck` crate to cast the slice of f32 to a &[u8] to be uploaded to the GPU. + let input_data_buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("Input"), + size: INPUT_BUF_SIZE, + usage: wgpu::BufferUsages::union(BufferUsages::STORAGE, BufferUsages::COPY_DST), + mapped_at_creation: false, + }); + + let data_size_buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("Data Size"), + size: SIZES_BUF_SIZE, + usage: wgpu::BufferUsages::union(BufferUsages::STORAGE, BufferUsages::COPY_DST), + mapped_at_creation: false, + }); + + // Now we create a buffer to store the output data. + let output_data_buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("Output"), + size: OUTPUT_BUF_SIZE, + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC, + mapped_at_creation: false, + }); + + // Finally we create a buffer which can be read by the CPU. This buffer is how we will read + // the data. We need to use a separate buffer because we need to have a usage of `MAP_READ`, + // and that usage can only be used with `COPY_DST`. + let download_buffer_0 = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("Download_0"), + size: OUTPUT_BUF_SIZE, + usage: wgpu::BufferUsages::COPY_DST | wgpu::BufferUsages::MAP_READ, + mapped_at_creation: false, + }); + + // We use a second buffer to flip stuff + let download_buffer_1 = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("Download_1"), + size: OUTPUT_BUF_SIZE, + usage: wgpu::BufferUsages::COPY_DST | wgpu::BufferUsages::MAP_READ, + mapped_at_creation: false, + }); + + let dl_bufs = (download_buffer_0, download_buffer_1); + + // A bind group layout describes the types of resources that a bind group can contain. Think + // of this like a C-style header declaration, ensuring both the pipeline and bind group agree + // on the types of resources. + let bind_group_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: None, + entries: &[ + // Input buffer + wgpu::BindGroupLayoutEntry { + binding: 0, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: true }, + // This is the size of a single element in the buffer. + min_binding_size: Some(NonZeroU64::new(512).unwrap()), + has_dynamic_offset: false, + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 1, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: true }, + // This is the size of a single element in the buffer. + min_binding_size: Some(NonZeroU64::new(4).unwrap()), + has_dynamic_offset: false, + }, + count: None, + }, + + // Output buffer + wgpu::BindGroupLayoutEntry { + binding: 2, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: false }, + // This is the size of a single element in the buffer. + min_binding_size: Some(NonZeroU64::new(32).unwrap()), + has_dynamic_offset: false, + }, + count: None, + }, + ], + }); + + // The bind group contains the actual resources to bind to the pipeline. + // + // Even when the buffers are individually dropped, wgpu will keep the bind group and buffers + // alive until the bind group itself is dropped. + let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor { + label: None, + layout: &bind_group_layout, + entries: &[ + wgpu::BindGroupEntry { + binding: 0, + resource: input_data_buffer.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 1, + resource: data_size_buffer.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 2, + resource: output_data_buffer.as_entire_binding(), + }, + ], + }); + + // The pipeline layout describes the bind groups that a pipeline expects + let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { + label: None, + bind_group_layouts: &[&bind_group_layout], + immediate_size: 0, + }); + + // The pipeline is the ready-to-go program state for the GPU. It contains the shader modules, + // the interfaces (bind group layouts) and the shader entry point. + let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: None, + layout: Some(&pipeline_layout), + module: &module, + entry_point: Some("main"), + compilation_options: wgpu::PipelineCompilationOptions::default(), + cache: None, + }); + + + let mut nibbles: [u16; OUTPUT_SIZE as usize/ 8] = [0;4]; + let mut b = bits as u16; + for n in 0..4 { + if b > 64 { + nibbles[n] = 64; + b -= 64; + } else { + nibbles[n] = b; + break; + } + } + + // == Begin repeating part + let mut found: Option<[u8;32]> = None; + let mut found_input: Option<String> = None; + let mut nonces_1: Vec<NonceElement>= Vec::with_capacity(BATCH_SIZE as usize); + let mut nonces_2: Vec<NonceElement>= Vec::with_capacity(BATCH_SIZE as usize); + let mut nonce_bufs = (nonces_1, nonces_2); + let mut prev_time: Instant = Instant::now(); + let mut prev_submission_index = None; + println!("starting!"); + for x in 0 .. MAX_ITER { + let data_upload_buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("Upload Data"), + size: INPUT_BUF_SIZE, + usage: wgpu::BufferUsages::COPY_SRC | wgpu::BufferUsages::MAP_WRITE, + mapped_at_creation: true, + }); + let size_upload_buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("Upload Size"), + size: SIZES_BUF_SIZE, + usage: wgpu::BufferUsages::COPY_SRC | wgpu::BufferUsages::MAP_WRITE, + mapped_at_creation: true, + }); + let mut input_slice = data_upload_buffer.get_mapped_range_mut(..); + let mut len_slice = size_upload_buffer.get_mapped_range_mut(..); + + let nonces = swaptuple_get_mut(&mut nonce_bufs, x); + for i in 0 .. BATCH_SIZE { + if let Some(element) = iter.next() { + let inputdata = element.to_hash_input(&value, &prev).into_bytes(); + let inputlen = inputdata.len() as u32; + let lenbytes = inputlen.to_le_bytes(); + + // println!("{}: {}, {}", i, inputlen*8, str::from_utf8(&inputdata).unwrap()); + input_slice[(i * INPUT_SIZE) as usize .. (i*INPUT_SIZE + inputlen as u64) as usize].copy_from_slice(&inputdata); + len_slice[(i * INSIZE_SIZE) as usize .. ((i+1) * INSIZE_SIZE) as usize].copy_from_slice(&lenbytes); + // println!("In data {:?}", &input_slice[(i * INPUT_SIZE) as usize .. (i*INPUT_SIZE + inputlen as u64) as usize]); + // println!("In length {:?}", &len_slice[(i * INSIZE_SIZE) as usize .. ((i+1) * INSIZE_SIZE) as usize]); + nonces.push(element); + } else { + panic!("Infinite iter died") + } + } + // println!("Hex slice:\n{}", hex::encode(&input_slice[..])); + drop(input_slice); // The command encoder allows us to record commands that we will later submit to the GPU. + drop(len_slice); + data_upload_buffer.unmap(); + size_upload_buffer.unmap(); + + let mut encoder = + device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None }); + + // We add a copy operation to the encoder. This will copy the data from the output buffer on the + // GPU to the download buffer on the CPU. + encoder.copy_buffer_to_buffer( + &data_upload_buffer, + 0, + &input_data_buffer, + 0, + data_upload_buffer.size(), + ); + encoder.copy_buffer_to_buffer( + &size_upload_buffer, + 0, + &data_size_buffer, + 0, + size_upload_buffer.size(), + ); + + // A compute pass is a single series of compute operations. While we are recording a compute + // pass, we cannot record to the encoder. + let mut compute_pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: None, + timestamp_writes: None, + }); + + // Set the pipeline that we want to use + compute_pass.set_pipeline(&pipeline); + // Set the bind group that we want to use + compute_pass.set_bind_group(0, &bind_group, &[]); + + // Now we dispatch a series of workgroups. Each workgroup is a 3D grid of individual programs. + // + // We defined the workgroup size in the shader as 64x1x1. So in order to process all of our + // inputs, we ceiling divide the number of inputs by 64. If the user passes 32 inputs, we will + // dispatch 1 workgroups. If the user passes 65 inputs, we will dispatch 2 workgroups, etc. + + let rest = BATCH_SIZE % 64; + let mut workgroup_count = BATCH_SIZE - rest / 64; + if rest > 0 { + workgroup_count += 1; + } + compute_pass.dispatch_workgroups(workgroup_count as u32, 1, 1); + + // Now we drop the compute pass, giving us access to the encoder again. + drop(compute_pass); + + // We add a copy operation to the encoder. This will copy the data from the output buffer on the + // GPU to the download buffer on the CPU. + encoder.copy_buffer_to_buffer( + &output_data_buffer, + 0, + swaptuple_get(&dl_bufs, x), + 0, + output_data_buffer.size(), + ); + + // We finish the encoder, giving us a fully recorded command buffer. + let command_buffer = encoder.finish(); + + + + let sub_index = queue.submit([command_buffer]); + // println!("A pass has been submitted"); + // Wait for the GPU to finish working on the submitted work. This doesn't work on WebGPU, so we would need + // to rely on the callback to know when the buffer is mapped. + + if x > 0 { + let nonces = swaptuple_get_mut(&mut nonce_bufs, x -1); + let download_buffer = swaptuple_get(&dl_bufs, x-1); + let buffer_slice = download_buffer.slice(..); + buffer_slice.map_async(wgpu::MapMode::Read, |_| {}); + device.poll(wgpu::PollType::Wait { submission_index: prev_submission_index, timeout: None }).unwrap(); + // In this case we know exactly when the mapping will be finished, + // so we don't need to do anything in the callback. + let data = buffer_slice.get_mapped_range(); + // println!("Out data {:?}", &data[..]); + // println!("Out hash {:?}", hex::encode(&data[..])); + let results: &[u8] = bytemuck::cast_slice(&data[..]); + + // println!("Full buffer: {}", hex::encode(results) ); + for i in 0..(BATCH_SIZE as usize) { + let result: &[u8;32] = &results[32*i..32*(i+1)].try_into().unwrap(); + // println!("Hash of {}, {}", hex::encode(nonces.get(i).unwrap().bytes.clone()), hex::encode(bytemuck::cast_slice(result))); + let mut correct = true; + for n in 0..4 { + if u64::from_be_bytes(result[n*8..(n+1)*8].try_into().unwrap()).leading_zeros() < nibbles[n] as u32{ + correct = false; + break; + } + } + if correct { + found = Some(*result); + let nonce = nonces.get(i).expect("We pushed all those nonces, right?").bytes.clone(); + found_input = Some(String::from_utf8(nonce).unwrap()); + // break; + } + } + if found.is_some() { + break; + } + drop(data); + download_buffer.unmap(); + nonces.clear(); + let now = Instant::now(); + let delta = (now - prev_time).as_secs_f64(); + let hashrate = BATCH_SIZE as f64 / delta; + println!("\x1b[A\rFinished batch {}, {} hashes processed ({:.2} H/s) (delta: {}) ", x + 1, x * BATCH_SIZE, hashrate, delta); + prev_time = Instant::now(); + } + prev_submission_index = Some(sub_index); + + } + + match found { + Some(f) => { + let hashbytes: &[u8] = bytemuck::cast_slice(&f); + let nonce = found_input.unwrap(); + println!("Found! {:?}", hex::encode(hashbytes)); + println!("Full block: {}{}{}", value, prev , nonce); + println!("Nonce: \"{}\"", nonce); + } + None => { + println!("Could not find a valid hash") + } + } + + + + println!("Done for now") + + +} diff --git a/src/shader.wgsl b/src/shader.wgsl new file mode 100644 index 0000000..39b53d5 --- /dev/null +++ b/src/shader.wgsl @@ -0,0 +1,208 @@ +struct SHA256_CTX { + data : array<u32, 64>, + datalen : u32, + bitlen : array<u32, 2>, + state : array<u32, 8>, + info : u32, +}; + +@group(0) @binding(0) var<storage, read> inputs : array<u32>; +@group(0) @binding(1) var<storage, read> inputSizes : array<u32>; +@group(0) @binding(2) var<storage, read_write> results : array<u32>; + +const SHA256_BLOCK_SIZE = 32; +const INPUT_BLOCK_SIZE = 64; +const INPUT_BLOCKS_MAX = 8; +const INPUT_SLOT_SIZE = INPUT_BLOCK_SIZE * INPUT_BLOCKS_MAX / 4; + +const k = array<u32, 64> ( + 0x428a2f98,0x71374491,0xb5c0fbcf,0xe9b5dba5,0x3956c25b,0x59f111f1,0x923f82a4,0xab1c5ed5, + 0xd807aa98,0x12835b01,0x243185be,0x550c7dc3,0x72be5d74,0x80deb1fe,0x9bdc06a7,0xc19bf174, + 0xe49b69c1,0xefbe4786,0x0fc19dc6,0x240ca1cc,0x2de92c6f,0x4a7484aa,0x5cb0a9dc,0x76f988da, + 0x983e5152,0xa831c66d,0xb00327c8,0xbf597fc7,0xc6e00bf3,0xd5a79147,0x06ca6351,0x14292967, + 0x27b70a85,0x2e1b2138,0x4d2c6dfc,0x53380d13,0x650a7354,0x766a0abb,0x81c2c92e,0x92722c85, + 0xa2bfe8a1,0xa81a664b,0xc24b8b70,0xc76c51a3,0xd192e819,0xd6990624,0xf40e3585,0x106aa070, + 0x19a4c116,0x1e376c08,0x2748774c,0x34b0bcb5,0x391c0cb3,0x4ed8aa4a,0x5b9cca4f,0x682e6ff3, + 0x748f82ee,0x78a5636f,0x84c87814,0x8cc70208,0x90befffa,0xa4506ceb,0xbef9a3f7,0xc67178f2 +); + +fn ROTLEFT(a : u32, b : u32) -> u32{return (((a) << (b)) | ((a) >> (32-(b))));} +fn ROTRIGHT(a : u32, b : u32) -> u32{return (((a) >> (b)) | ((a) << (32-(b))));} + +fn CH(x : u32, y : u32, z : u32) -> u32{return (((x) & (y)) ^ (~(x) & (z)));} +fn MAJ(x : u32, y : u32, z : u32) -> u32{return (((x) & (y)) ^ ((x) & (z)) ^ ((y) & (z)));} +fn EP0(x : u32) -> u32{return (ROTRIGHT(x,2) ^ ROTRIGHT(x,13) ^ ROTRIGHT(x,22));} +fn EP1(x : u32) -> u32{return (ROTRIGHT(x,6) ^ ROTRIGHT(x,11) ^ ROTRIGHT(x,25));} +fn SIG0(x : u32) -> u32{return (ROTRIGHT(x,7) ^ ROTRIGHT(x,18) ^ ((x) >> 3));} +fn SIG1(x : u32) -> u32{return (ROTRIGHT(x,17) ^ ROTRIGHT(x,19) ^ ((x) >> 10));} + +fn sha256_transform(ctx : ptr<function, SHA256_CTX>) +{ + var a : u32; + var b : u32; + var c : u32; + var d : u32; + var e : u32; + var f : u32; + var g : u32; + var h : u32; + var i : u32 = 0; + var j : u32 = 0; + var t1 : u32; + var t2 : u32; + var m : array<u32, 64> ; + + + while(i < 16) { + m[i] = ((*ctx).data[j] << 24) | ((*ctx).data[j + 1] << 16) | ((*ctx).data[j + 2] << 8) | ((*ctx).data[j + 3]); + i++; + j += 4; + } + + while(i < 64) { + m[i] = SIG1(m[i - 2]) + m[i - 7] + SIG0(m[i - 15]) + m[i - 16]; + i++; + } + + a = (*ctx).state[0]; + b = (*ctx).state[1]; + c = (*ctx).state[2]; + d = (*ctx).state[3]; + e = (*ctx).state[4]; + f = (*ctx).state[5]; + g = (*ctx).state[6]; + h = (*ctx).state[7]; + + i = 0; + for (; i < 64; i++) { + t1 = h + EP1(e) + CH(e,f,g) + k[i] + m[i]; + t2 = EP0(a) + MAJ(a,b,c); + h = g; + g = f; + f = e; + e = d + t1; + d = c; + c = b; + b = a; + a = t1 + t2; + } + + + (*ctx).state[0] += a; + (*ctx).state[1] += b; + (*ctx).state[2] += c; + (*ctx).state[3] += d; + (*ctx).state[4] += e; + (*ctx).state[5] += f; + (*ctx).state[6] += g; + (*ctx).state[7] += h; +} + + +fn sha256_update(ctx : ptr<function, SHA256_CTX>, len : u32, index: u32) +{ + for (var i :u32 = (INPUT_SLOT_SIZE * index); i < ((INPUT_SLOT_SIZE*index)+len); i++) { + (*ctx).data[(*ctx).datalen] = inputs[i]; + (*ctx).datalen++; + if ((*ctx).datalen == 64) { + sha256_transform(ctx); + + if ((*ctx).bitlen[0] > 0xffffffff - (512)){ + (*ctx).bitlen[1]++; + } + (*ctx).bitlen[0] += 512; + + + (*ctx).datalen = 0; + } + } +} + +fn sha256_final(ctx : ptr<function, SHA256_CTX>, hash: ptr<function, array<u32, 8>> ) +{ + var i : u32 = (*ctx).datalen; + + if ((*ctx).datalen < 56) { + (*ctx).data[i] = 0x80; + i++; + while (i < 56){ + (*ctx).data[i] = 0x00; + i++; + } + } + else { + (*ctx).data[i] = 0x80; + i++; + while (i < 64){ + (*ctx).data[i] = 0x00; + i++; + } + sha256_transform(ctx); + for (var i = 0; i < 56 ; i++) { + (*ctx).data[i] = 0; + } + } + + if ((*ctx).bitlen[0] > 0xffffffff - (*ctx).datalen * 8) { + (*ctx).bitlen[1]++; + } + (*ctx).bitlen[0] += (*ctx).datalen * 8; + + + (*ctx).data[63] = (*ctx).bitlen[0]; + (*ctx).data[62] = (*ctx).bitlen[0] >> 8; + (*ctx).data[61] = (*ctx).bitlen[0] >> 16; + (*ctx).data[60] = (*ctx).bitlen[0] >> 24; + (*ctx).data[59] = (*ctx).bitlen[1]; + (*ctx).data[58] = (*ctx).bitlen[1] >> 8; + (*ctx).data[57] = (*ctx).bitlen[1] >> 16; + (*ctx).data[56] = (*ctx).bitlen[1] >> 24; + sha256_transform(ctx); + + for (i = 0; i < 8; i++) { + (*hash)[i] = (*ctx).state[i]; + + } + // for (i = 0; i < 4; i++) { + // (*hash)[i] = ((*ctx).state[0] >> (24 - i * 8)) & 0x000000ff; + // (*hash)[i + 4] = ((*ctx).state[1] >> (24 - i * 8)) & 0x000000ff; + // (*hash)[i + 8] = ((*ctx).state[2] >> (24 - i * 8)) & 0x000000ff; + // (*hash)[i + 12] = ((*ctx).state[3] >> (24 - i * 8)) & 0x000000ff; + // (*hash)[i + 16] = ((*ctx).state[4] >> (24 - i * 8)) & 0x000000ff; + // (*hash)[i + 20] = ((*ctx).state[5] >> (24 - i * 8)) & 0x000000ff; + // (*hash)[i + 24] = ((*ctx).state[6] >> (24 - i * 8)) & 0x000000ff; + // (*hash)[i + 28] = ((*ctx).state[7] >> (24 - i * 8)) & 0x000000ff; + // } +} + +@compute @workgroup_size(1, 1,1) +fn main(@builtin(global_invocation_id) global_id : vec3<u32>) { + var index = global_id.x; + let array_length = arrayLength(&inputSizes); + if (index >= array_length) { + return; + } + var ctx : SHA256_CTX; + var buf : array<u32, 8>; + + // CTX INIT + ctx.datalen = 0; + ctx.bitlen[0] = 0; + ctx.bitlen[1] = 0; + ctx.state[0] = 0x6a09e667; + ctx.state[1] = 0xbb67ae85; + ctx.state[2] = 0x3c6ef372; + ctx.state[3] = 0xa54ff53a; + ctx.state[4] = 0x510e527f; + ctx.state[5] = 0x9b05688c; + ctx.state[6] = 0x1f83d9ab; + ctx.state[7] = 0x5be0cd19; + + sha256_update(&ctx, inputSizes[index], index); + sha256_final(&ctx, &buf); + + + for (var i=(index * SHA256_BLOCK_SIZE); i < ((index+1) * SHA256_BLOCK_SIZE); i++) { + results[i] = buf[i]; + } +} diff --git a/src/shader_own.wgsl b/src/shader_own.wgsl new file mode 100644 index 0000000..90e1a84 --- /dev/null +++ b/src/shader_own.wgsl @@ -0,0 +1,226 @@ +struct SHA256_CTX { + data : array<u32, 64>, + datalen: u32, + bitlen: array<u32, 2>, + state: array<u32, 8>, + info: array<u32, 8>, +} + +const SHA256_BLOCK_SIZE = 32; +const INPUT_BLOCK_SIZE = 64; +const INPUT_BLOCK_MAX = 8; +const INPUT_PACKED_SIZE = (INPUT_BLOCK_SIZE * INPUT_BLOCK_MAX / 4); + +struct InputBlocks { + data: array<u32, 128>, +} +struct OutputHash { + data: array<u32, 8>, +} + +@group(0) @binding(0) var<storage, read> inputs: array<InputBlocks>; +@group(0) @binding(1) var<storage, read> sizes: array<u32>; +@group(0) @binding(2) var<storage, read_write> results: array<OutputHash>; + + +const k = array<u32, 64> ( + 0x428a2f98,0x71374491,0xb5c0fbcf,0xe9b5dba5,0x3956c25b,0x59f111f1,0x923f82a4,0xab1c5ed5, + 0xd807aa98,0x12835b01,0x243185be,0x550c7dc3,0x72be5d74,0x80deb1fe,0x9bdc06a7,0xc19bf174, + 0xe49b69c1,0xefbe4786,0x0fc19dc6,0x240ca1cc,0x2de92c6f,0x4a7484aa,0x5cb0a9dc,0x76f988da, + 0x983e5152,0xa831c66d,0xb00327c8,0xbf597fc7,0xc6e00bf3,0xd5a79147,0x06ca6351,0x14292967, + 0x27b70a85,0x2e1b2138,0x4d2c6dfc,0x53380d13,0x650a7354,0x766a0abb,0x81c2c92e,0x92722c85, + 0xa2bfe8a1,0xa81a664b,0xc24b8b70,0xc76c51a3,0xd192e819,0xd6990624,0xf40e3585,0x106aa070, + 0x19a4c116,0x1e376c08,0x2748774c,0x34b0bcb5,0x391c0cb3,0x4ed8aa4a,0x5b9cca4f,0x682e6ff3, + 0x748f82ee,0x78a5636f,0x84c87814,0x8cc70208,0x90befffa,0xa4506ceb,0xbef9a3f7,0xc67178f2 +); + +const masks = array<u32, 4> ( + 0x000000FFu, + 0x0000FF00u, + 0x00FF0000u, + 0xFF000000u, +); + +const full_masks = array<u32, 5> ( + 0x00000000u, + 0x000000FFu, + 0x0000FFFFu, + 0x00FFFFFFu, + 0xFFFFFFFFu, +); + +fn ROTLEFT(a : u32, b : u32) -> u32{return (((a) << (b)) | ((a) >> (32-(b))));} +fn ROTRIGHT(a : u32, b : u32) -> u32{return (((a) >> (b)) | ((a) << (32-(b))));} + +fn CH(e : u32, f : u32, g : u32) -> u32{return (((e) & (f)) ^ (~(e) & (g)));} +fn MAJ(a : u32, b : u32, c: u32) -> u32{return (((a) & (b)) ^ ((a) & (c)) ^ ((b) & (c)));} +fn EP0(x : u32) -> u32{return (ROTRIGHT(x,2) ^ ROTRIGHT(x,13) ^ ROTRIGHT(x,22));} +fn EP1(x : u32) -> u32{return (ROTRIGHT(x,6) ^ ROTRIGHT(x,11) ^ ROTRIGHT(x,25));} +fn SIG0(x : u32) -> u32{return (ROTRIGHT(x,7) ^ ROTRIGHT(x,18) ^ ((x) >> 3));} +fn SIG1(x : u32) -> u32{return (ROTRIGHT(x,17) ^ ROTRIGHT(x,19) ^ ((x) >> 10));} + +fn flip_endian(x: u32) -> u32 { + var y: u32 = 0; + // 11 22 33 44 => 44 33 22 11 + // 11 >> 3*8 << 0*8; + // 22 >> 2*8 << 1*8; + // 33 >> 1*8 << 2*8; + // 44 >> 0*8 << 3*8; + for (var i: u32; i < 4; i++) { + y |= ((x >> (i*8)) & 0xFF) << ((3-i) * 8); + } + return y; +} + +fn sha256_init(ctx: ptr<function, SHA256_CTX>) { + + // CTX INIT + (*ctx).datalen = 0; + (*ctx).bitlen[0] = 0; + (*ctx).bitlen[1] = 0; + (*ctx).state[0] = 0x6a09e667; + (*ctx).state[1] = 0xbb67ae85; + (*ctx).state[2] = 0x3c6ef372; + (*ctx).state[3] = 0xa54ff53a; + (*ctx).state[4] = 0x510e527f; + (*ctx).state[5] = 0x9b05688c; + (*ctx).state[6] = 0x1f83d9ab; + (*ctx).state[7] = 0x5be0cd19; +} + +fn sha256_transform(ctx: ptr<function, SHA256_CTX>) { + var a : u32 = (*ctx).state[0]; + var b : u32 = (*ctx).state[1]; + var c : u32 = (*ctx).state[2]; + var d : u32 = (*ctx).state[3]; + var e : u32 = (*ctx).state[4]; + var f : u32 = (*ctx).state[5]; + var g : u32 = (*ctx).state[6]; + var h : u32 = (*ctx).state[7]; + var t1 : u32; + var t2 : u32; + var m: array<u32,64>; + + + for (var x: u32 = 0; x < 16; x++) { + m[x] = flip_endian((*ctx).data[x]); + } + for (var x: u32 = 16; x < 64; x++) { + // W_16 = w_0 + sig0(w_1) + w_9 + sig1(w_14) + m[x]= m[x-16] + SIG0(m[x-15]) + m[x-7] + SIG1(m[x-2]); + // if (x >= 16 && x < 24) { + // var z = x - 16; + // (*ctx).info[z] = m[x]; + // } + } + + + for (var x: u32 = 0; x < 64; x++) { + t1 = h + EP1(e) + CH(e, f, g) + k[x] + m[x]; + t2 = EP0(a) + MAJ(a,b,c); + h = g; + g = f; + f = e; + e = d + t1; + d = c; + c = b; + b = a; + a = t1 + t2; + + } + + (*ctx).state[0] += a; + (*ctx).state[1] += b; + (*ctx).state[2] += c; + (*ctx).state[3] += d; + (*ctx).state[4] += e; + (*ctx).state[5] += f; + (*ctx).state[6] += g; + (*ctx).state[7] += h; + + for (var x: u32 = 0; x < 64; x++) { + (*ctx).data[x] = m[x]; + } + + +} + +fn sha256_append_len(ctx: ptr<function, SHA256_CTX>, len: u32) { + if ((*ctx).bitlen[0] > 0xffffffff - (len)) { + (*ctx).bitlen[1]++; + } + (*ctx).bitlen[0] += len; +} + + +fn sha256_update(ctx: ptr<function, SHA256_CTX>, index: u32) { + var len = sizes[index]; + var block_index = 0u; + while len > 64 { + for (var i : u32 = 0; i < 16; i++) { + (*ctx).data[i] = inputs[index].data[16 * block_index + i]; + } + sha256_transform(ctx); + sha256_append_len(ctx, 512); + block_index += 1; + len -= 64; + } + var imax = len/4; + if (len % 4 > 0) {imax += 1;} + for (var i : u32 = 0; i < (imax); i++) { + (*ctx).data[i] = inputs[index].data[16 * block_index + i]; + } + for(var i = imax; i < 16; i ++) { + (*ctx).data[i] = 0; + } + + var offset : u32 = (len) % 4; + var last_word: u32 = (len - offset) / 4; + // if (offset == 0) {last_word += 1;} + var padding_word: u32 = (0x80u << (8 * offset)); + (*ctx).data[last_word] = ((*ctx).data[last_word] & full_masks[offset]) | padding_word; + + for (var i: u32 = 0; i < 8; i++) { + (*ctx).info[i] = (*ctx).data[i]; + } + for (var i: u32 = 14; i < 16; i++) { + (*ctx).info[i] = (*ctx).data[i]; + } + + if (len > 56) { + sha256_transform(ctx); + for(var i = 0; i < 14; i ++) { + (*ctx).data[i] = 0; + } + } + + var final_len: u32 = (len*8); + sha256_append_len(ctx, final_len); + var upper = flip_endian((*ctx).bitlen[0]); + var lower = flip_endian((*ctx).bitlen[1]); + + (*ctx).data[15] = upper; + (*ctx).data[14] = lower; + + sha256_transform(ctx); +} + +@compute @workgroup_size(64,1,1) +fn main(@builtin(global_invocation_id) global_id: vec3<u32>) { + var ctx: SHA256_CTX; + var index: u32 = global_id.x; + sha256_init(&ctx); + sha256_update(&ctx, index); + + for(var i: u32 = 0; i < 8; i++) { + results[index].data[i] = flip_endian(ctx.state[i]); + } + // for(var i: u32 = 0; i < 8; i++) { + // results[index].data[i] = inputs[index].data[i]; + // } + + // results[index].data[0] = ctx.info; +} + + + |