ソースを参照

Cargo check works, still hasnt run, majority of compute pipeline finished

Ghastrod 1 年間 前
コミット
1fdf07e46c
4 ファイル変更159 行追加29 行削除
  1. 80 27
      src/compute/compute_pipeline.rs
  2. 1 1
      src/compute/mod.rs
  3. 14 1
      src/lib.rs
  4. 64 0
      src/shaders/Heightmap_compute.wgsl

+ 80 - 27
src/compute/compute_pipeline.rs

@@ -1,38 +1,78 @@
-use wgpu::{include_wgsl, util::DeviceExt, BufferUsages, ComputePipeline};
+use serde::de;
+use wgpu::{include_wgsl, util::DeviceExt, BindGroup, BufferUsages, ComputePipeline};
 
-pub fn new_compute_pipeline(device: &wgpu::Device)-> ComputePipeline{
+pub const SIZE:u32 = 256;
+pub fn new_compute_pipeline(device: &wgpu::Device)-> (ComputePipeline,BindGroup,(u32,u32)){
 
-    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 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/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_shader = device.create_shader_module(include_wgsl!("../shaders/Heightmap_compute.wgsl"));
+
+
+    let output_texture_size = wgpu::Extent3d{
+        width: SIZE,
+        height: SIZE,
+        depth_or_array_layers: 0
+    };
+    let output_texture = device.create_texture(&wgpu::TextureDescriptor{
+        label: Some("Output Heightmap Texture"),
+        size: output_texture_size,
+        mip_level_count: 1,
+        sample_count: 1,
+        dimension: wgpu::TextureDimension::D2,
+        format: wgpu::TextureFormat::Rgba8UnormSrgb,
+        usage: wgpu::TextureUsages::STORAGE_BINDING | wgpu::TextureUsages::COPY_DST | wgpu::TextureUsages::COPY_SRC,
+        view_formats: &[]
     });
 
-    // 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 output_texture_view = output_texture.create_view(&wgpu::TextureViewDescriptor::default());
+
 
 
+    let heightmap_texture_bind_group_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor{
+        label: Some("Heightmap_texture_bind_group_layout"),
+        entries: &[
+            wgpu::BindGroupLayoutEntry{
+                binding: 0,
+                visibility: wgpu::ShaderStages::COMPUTE | wgpu::ShaderStages::VERTEX,
+                ty: wgpu::BindingType::StorageTexture{
+                    access: wgpu::StorageTextureAccess::ReadWrite,
+                    format: wgpu::TextureFormat::Rgba8UnormSrgb,
+                    view_dimension: wgpu::TextureViewDimension::D2
+                },
+                count: None
+            }
+        ]
+    });
+
+    let heightmap_texture_bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor{
+        label: Some("Heightmap texture bind group"),
+        layout: &heightmap_texture_bind_group_layout,
+        entries: &[
+            wgpu::BindGroupEntry{
+                binding: 0,
+                resource: wgpu::BindingResource::TextureView(&output_texture_view)
+            }
+        ]
+    });
+
+    //Pipeline Stuff
+    let compute_pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor{
+        label: Some("Compute Pipeline Layout 1"),
+        bind_group_layouts: &[
+            &heightmap_texture_bind_group_layout
+        ],
+        push_constant_ranges: &[]
+    });
     let compute_pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor{
         label: Some("Compute Pipeline 1"),
         layout: Some(&compute_pipeline_layout),
@@ -40,5 +80,18 @@ pub fn new_compute_pipeline(device: &wgpu::Device)-> ComputePipeline{
         entry_point: "@compute"
     });
 
-    return compute_pipeline
-}
+    let dispatch = compute_work_group_count((SIZE,SIZE), (16,16));
+    return (compute_pipeline,heightmap_texture_bind_group,dispatch)
+}
+
+pub fn compute_work_group_count(
+    (width, height): (u32, u32),
+    (workgroup_width, workgroup_height): (u32, u32),
+) -> (u32, u32) {
+    let x = (width + workgroup_width - 1) / workgroup_width;
+    let y = (height + workgroup_height - 1) / workgroup_height;
+
+    (x, y)
+}
+// compute_work_group_count((256,256),(16,16)) renvoie (16,16)
+// et compute_work_group_count((512,512),(16,16)) renvoie (32,32)

+ 1 - 1
src/compute/mod.rs

@@ -1 +1 @@
-mod compute_pipeline;
+pub mod compute_pipeline;

+ 14 - 1
src/lib.rs

@@ -1,7 +1,7 @@
 use std::{iter, time::{Duration, Instant}};
 
 use camera::{camera_buffer, camera_controller, camera_struct::Camera, camera_uniform};
-use wgpu::{include_wgsl, util::DeviceExt, BlendState, ColorTargetState, PipelineLayoutDescriptor, RenderPipeline, RenderPipelineDescriptor};
+use wgpu::{include_wgsl, util::DeviceExt, BlendState, ColorTargetState, ComputePipeline, PipelineLayoutDescriptor, RenderPipeline, RenderPipelineDescriptor};
 use winit::{
     event::*,
     event_loop::EventLoop,
@@ -13,6 +13,7 @@ mod object;
 mod sky;
 mod wireframe;
 mod texture;
+mod compute;
 
 #[cfg(target_arch = "wasm32")]
 use wasm_bindgen::prelude::*;
@@ -48,6 +49,9 @@ struct State<'a> {
     //Heightmap Stuff
     texture_bind_group: wgpu::BindGroup,
     diffuse_texture: texture::texture_struct::Texture,
+    compute_pipeline: ComputePipeline,
+    heightmap2_bindgroup: wgpu::BindGroup,
+    dispatch: (u32,u32),
 }
 
 impl<'a> State<'a> {
@@ -233,6 +237,9 @@ impl<'a> State<'a> {
 
         let wireframe_pipeline = wireframe::wireframe_pipeline::new(&device, &config, &camera_bind_group_layout, &diffuse_bind_group_layout);
         let wireframe = true;
+
+        let x = compute::compute_pipeline::new_compute_pipeline(&device);
+        //Compute
         Self {
             wireframe,
             surface,
@@ -262,6 +269,9 @@ impl<'a> State<'a> {
             //Heightmap Stuff
             diffuse_texture: diffuse_texture,
             texture_bind_group: diffuse_bind_group,
+            compute_pipeline: x.0,
+            heightmap2_bindgroup: x.1,
+            dispatch: x.2
         }
     }
 
@@ -336,6 +346,9 @@ impl<'a> State<'a> {
                 occlusion_query_set: None,
                 timestamp_writes: None,
             });
+            //Compute pipeline
+
+            //Render Pipeline
             //render_pass.set_pipeline(&self.sky_pipeline);
 
             if self.wireframe {

+ 64 - 0
src/shaders/Heightmap_compute.wgsl

@@ -0,0 +1,64 @@
+//Constants
+let PI: f32 = 3.14159265358979323846;
+let GRAVITY:f32= 9.81;
+
+//Structs
+struct Complex{
+    a: f32,
+    b: f32
+}
+
+fn add_complex(x:Complex, y:Complex)-> Complex{
+    return Complex(
+        x.a + y.a,
+        x.b + y.b 
+    )
+}
+
+fn multiply_complex(x:Complex, y: Complex)->Complex{
+    return Complex(
+        (x.a * y.a) - (x.b * y.b),
+        (y.b * x.a) + (y.a * x.a)
+    )
+}
+fn angle_to_complex(x:f32)-> Complex{
+    return Complex(
+        cos(x),
+        sin(x)
+    )
+}
+
+struct ComputeInput{
+    @builtin(local_invocation_id) local_invocation_id: vec3<u32>,
+    @builtin(local_invocation_index) local_invocation_index: u32,
+    @builtin(global_invocation_index) global_invocation_index: vec3<u32>,
+    @builtin(workgroup_id) workgroup_id: vec3<u32>,
+    @builtin(num_workgroups) num_workgroups: vec3<u32>
+}
+
+//Actual code
+
+@group(0) @binding(0) var heightmap_texture: texture_2d<rgba8unorm,write>;
+fn gaussian(x: f32, y:f32)-> f32{
+    let mean = 0.0;
+    let spread = 1.0;
+    let actual_x = x - mean;
+    let actual_y = y - mean;
+    let sigma_squared = spread * spread;
+    return (1 / (sqrt(2 * PI)* spread)) * exp(-((actual_x * actual_x) + (actual_y * actual_y)) / (2 * sigmaSqu));
+}
+
+@compute @workgroup_size(1,1)
+fn compute_main(in: ComputeInput){
+    let coordonnees = in.global_invocation_index.xy;
+    let dimensions = textureDimensions(heightmap_texture);
+
+    let pixel_color_1_channel = sin(coordonnees.x + coordonnees.y);
+
+    let vec4_pixel_color = vec4<f32>(pixel_color_1_channel);
+
+    textureStore(heightmap_texture, coordonnees, vec4_pixel_color)
+}
+
+
+