Restart of compute pipeline, after grayscale

src/compute/compute_pipeline.rs

@@ -1,14 +1,37 @@
-use wgpu::{include_wgsl, ComputePipeline};
+use wgpu::{include_wgsl, util::DeviceExt, BufferUsages, ComputePipeline};
 
 pub fn new_compute_pipeline(device: &wgpu::Device)-> ComputePipeline{
 
+    let u32_size = std::mem::size_of::<u64>() as wgpu::BufferAddress;
+    let time_buffer = device.create_buffer(&wgpu::BufferDescriptor{
+        label: Some("Time uniform"),
+        size: u32_size,
+        usage: BufferUsages::COPY_DST | BufferUsages::UNIFORM,
+        mapped_at_creation: false
+    });
+
+
     let compute_pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor{
         label: Some("Compute Pipeline Layout 1"),
         bind_group_layouts: &[],
         push_constant_ranges: &[]
     });
 
-    let compute_shader = device.create_shader_module(include_wgsl!("../shaders/compute.wgsl"));
+    let compute_shader = device.create_shader_module(include_wgsl!("../shaders/CS_FFTHorizontal.wgsl"));
+
+    let attente_buffer = device.create_buffer(&wgpu::BufferDescriptor{
+        label: Some("Buffer attente, qui reste sur le CPU"),
+        size: todo!(),
+        usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST,
+        mapped_at_creation: false,
+    });
+
+    let compute_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor{
+        label: Some("Compute Buffer"),
+        contents: bytemuck::cast_slice(todo!()),
+        usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC
+    });
+
 
     let compute_pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor{
         label: Some("Compute Pipeline 1"),

src/shaders/CS_FFTHorizontal.wgsl

@@ -0,0 +1,57 @@
+@group(0) @binding(0)
+var<uniform> u_input: texture_2d<f32>; // readonly image2D
+
+@group(0) @binding(1)
+var<storage, write> u_output: texture_2d<f32>; // writeonly image2D
+
+let PI: f32 = 3.14159265358979323846;
+
+struct Complex {
+  x: f32,
+  y: f32,
+};
+
+fn multiply_complex(a: Complex, b: Complex) -> Complex {
+  return Complex(
+    a.x * b.x - a.y * b.y,
+    a.y * b.x + a.x * b.y
+  );
+}
+
+fn add_complex(a: Complex, b: Complex) -> Complex{
+  return Complex(
+    a.x + b.x,
+    a.y + b.y
+  )
+}
+
+fn butterfly_operation(a: Complex, b: Complex, twiddle: Complex) -> vec4<f32> {
+  let twiddle_b: Complex = multiply_complex(twiddle, b);
+  let result: vec4<f32> = vec4<f32>(a.x + twiddle_b.x, a.y + twiddle_b.y, a.x - twiddle_b.x, a.y - twiddle_b.y);
+  return result;
+}
+
+@compute @workgroup_size(256)
+fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
+  let pixel_coord: vec2<u32> = vec2<u32>(global_id.x, global_id.y);
+
+  let thread_count: u32 = u32(u_total_count >> 1); // Integer division by 2
+  let thread_idx: u32 = pixel_coord.x;
+
+  let in_idx: u32 = thread_idx & (u_subseq_count - 1u);
+  let out_idx: u32 = ((thread_idx - in_idx) << 1) + in_idx;
+
+  let angle: f32 = -PI * (f32(in_idx) / f32(u_subseq_count));
+  let twiddle: Complex = Complex(cos(angle), sin(angle));
+
+  let a: vec4<f32> = textureSample(u_input, pixel_coord);
+  let b: vec4<f32> = textureSample(u_input, vec2<u32>(pixel_coord.x + thread_count, pixel_coord.y));
+
+  // Transforming two complex sequences independently and simultaneously
+
+  let result0: vec4<f32> = butterfly_operation(Complex(a.x, a.y), Complex(b.x, b.y), twiddle);
+  let result1: vec4<f32> = butterfly_operation(Complex(a.z, a.w), Complex(b.z, b.w), twiddle);
+
+  textureStore(u_output, vec2<u32>(out_idx, pixel_coord.y), vec4<f32>(result0.xy, result1.xy));
+  textureStore(u_output, vec2<u32>(out_idx + u_subseq_count, pixel_coord.y), vec4<f32>(result0.zw, result1.zw));
+}

src/shaders/CS_FFTVertical.wgsl

@@ -0,0 +1,56 @@
+@group(0) @binding(0)
+var<uniform> u_input: texture_2d<f32>; // readonly image2D
+
+@group(0) @binding(1)
+var<storage, write> u_output: texture_2d<f32>; // writeonly image2D
+
+let PI: f32 = 3.14159265358979323846;
+
+struct Complex {
+  x: f32,
+  y: f32,
+};
+
+fn multiply_complex(a: Complex, b: Complex) -> Complex {
+  return Complex(
+    a.x * b.x - a.y * b.y,
+    a.y * b.x + a.x * b.y
+  );
+}
+fn add_complex(a: Complex, b: Complex) -> Complex{
+  return Complex(
+    a.x + b.x,
+    a.y + b.y
+  )
+}
+
+fn butterfly_operation(a: Complex, b: Complex, twiddle: Complex) -> vec4<f32> {
+  let twiddle_b: Complex = multiply_complex(twiddle, b);
+  let result: vec4<f32> = vec4<f32>(a.x + twiddle_b.x, a.y + twiddle_b.y, a.x - twiddle_b.x, a.y - twiddle_b.y);
+  return result;
+}
+
+@compute @workgroup_size(256)
+fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
+  let pixel_coord: vec2<u32> = vec2<u32>(global_id.x, global_id.y);
+
+  let thread_count: u32 = u32(u_total_count >> 1); // Integer division by 2
+  let thread_idx: u32 = pixel_coord.y;
+
+  let in_idx: u32 = thread_idx & (u_subseq_count - 1u);
+  let out_idx: u32 = ((thread_idx - in_idx) << 1) + in_idx;
+
+  let angle: f32 = -PI * (f32(in_idx) / f32(u_subseq_count));
+  let twiddle: Complex = Complex(cos(angle), sin(angle));
+
+  let a: vec4<f32> = textureSample(u_input, pixel_coord);
+  let b: vec4<f32> = textureSample(u_input, vec2<u32>(pixel_coord.x, pixel_coord.y + thread_count));
+
+  // Transforming two complex sequences independently and simultaneously
+
+  let result0: vec4<f32> = butterfly_operation(Complex(a.x, a.y), Complex(b.x, b.y), twiddle);
+  let result1: vec4<f32> = butterfly_operation(Complex(a.z, a.w), Complex(b.z, b.w), twiddle);
+
+  textureStore(u_output, vec2<u32>(pixel_coord.x, out_idx), vec4<f32>(result0.xy, result1.xy));
+  textureStore(u_output, vec2<u32>(pixel_coord.x, out_idx + u_subseq_count), vec4<f32>(result0.zw, result1.zw));
+}

src/shaders/CS_mine.wgsl

@@ -0,0 +1,33 @@
+struct Complex{
+    x: f32,
+    y:f32
+};
+let PI: f32 = 3.14159265358979323846;
+
+fn add_complex(a:Complex, b:Complex)-> Complex{
+    return Complex(
+        a.x + b.x,
+        a.y + b.y
+    )
+}
+
+fn multiply_complex(a:Complex, b:Complex)-> Complex{
+    return Complex(
+        (a.x * b.x) - (a.y * b.y),
+        (a.x * b.y) + (b.x * a.y)
+    )
+}
+
+fn gaussian(x: f32, y:f32)-> f32{
+    //Mean = 0
+    let spread = 1.0;
+    let sigma_squared = spread * spread;
+    return (1 / (sqrt(2 * PI)* spread)) * exp(-((x * x) + (y * y)) / (2 * sigmaSqu));
+}
+
+fn angle_to_complex(x:f32)-> Complex{
+    return Complex(
+        cos(x),
+        sin(x)
+    )
+}

src/shaders/compute.wgsl

@@ -1,5 +0,0 @@
-@compute
-fn compute_main(){}
-
-@vertex
-fn debug(){}