working version

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;
+}
+
+
+