From e73a4cd05e429ca3e0c57db4010d54817bcfce0a Mon Sep 17 00:00:00 2001 From: Jarrod Doyle Date: Thu, 9 May 2024 15:19:46 +0100 Subject: [PATCH] Add raygen shader --- res/common/raygen.wgsl | 43 ++++++++++++++++++++++++++ src/main.rs | 30 ++++++++++++++++++- src/renderer.rs | 68 ++++++++++++++++++++++++++++++++++++++++-- 3 files changed, 137 insertions(+), 4 deletions(-) create mode 100644 res/common/raygen.wgsl diff --git a/res/common/raygen.wgsl b/res/common/raygen.wgsl new file mode 100644 index 0000000..0e6a631 --- /dev/null +++ b/res/common/raygen.wgsl @@ -0,0 +1,43 @@ +@group(0) @binding(0) var output: texture_storage_2d; +@group(0) @binding(1) var camera: Camera; +@group(0) @binding(2) var rays: array; + +struct Camera { + projection: mat4x4, + view: mat4x4, + pos: vec3, +}; + +struct Ray { + pos: vec3, + dist: f32, + dir: vec3, + hit: u32, +} + +@compute @workgroup_size(8, 8, 1) +fn compute(@builtin(global_invocation_id) global_id: vec3) { + let img_coord = global_id.xy; + let img_dims = textureDimensions(output); + + // This discards the extra pixels in cases where the image size isn't perfectly divisible by the kernel.xy + if (img_coord.x >= img_dims.x || img_coord.y >= img_dims.y) { + return; + } + + // Construct ray + let img_coord_frac = vec2(img_coord) / vec2(img_dims); + let screen_pos = img_coord_frac * 2.0 - vec2(1.0); + var ray_eye = camera.projection * vec4(screen_pos, -1.0, 0.0); + ray_eye = vec4(ray_eye.xy, -1.0, 0.0); + let ray_dir = normalize((camera.view * ray_eye).xyz); + let ray_pos = camera.pos; + + let idx = img_coord.x + img_coord.y * img_dims.x; + let ray = Ray(ray_pos, 0.0, ray_dir, 0); + rays[idx] = ray; + + // Temp + var color = vec4(f32(img_coord.x) / f32(img_dims.x), f32(img_coord.y) / f32(img_dims.y), 0.5, 1.0); + textureStore(output, img_coord, color); +} \ No newline at end of file diff --git a/src/main.rs b/src/main.rs index 0318612..c4c471e 100644 --- a/src/main.rs +++ b/src/main.rs @@ -1,3 +1,4 @@ +mod camera; mod renderer; use anyhow::Result; @@ -12,7 +13,30 @@ fn main() -> Result<()> { .with_vsync(false) .build(), )?; - let renderer = VoxelRenderer::new(&context)?; + + let mut camera_controller = camera::CameraController::new( + &context, + camera::Camera::new( + glam::Vec3 { + x: 4.01, + y: 4.01, + z: 20.0, + }, + -90.0_f32.to_radians(), + 0.0_f32.to_radians(), + ), + camera::Projection::new( + context.size.width, + context.size.height, + 90.0_f32.to_radians(), + 0.01, + 100.0, + ), + 10.0, + 0.25, + ); + + let renderer = VoxelRenderer::new(&context, &camera_controller)?; let passes: Vec> = vec![Box::new(renderer)]; event_loop.run(|event, elwt| match event { @@ -21,6 +45,10 @@ fn main() -> Result<()> { return; } + if camera_controller.process_events(&event) { + return; + } + if let WindowEvent::RedrawRequested = event { if let Err(err) = context.render(passes.as_slice()) { log::error!("{err:?}"); diff --git a/src/renderer.rs b/src/renderer.rs index 5603778..b0e764f 100644 --- a/src/renderer.rs +++ b/src/renderer.rs @@ -1,14 +1,27 @@ use anyhow::Result; use crawl::wgpu; +use crate::camera::CameraController; + +#[repr(C)] +#[derive(Debug, Default, Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)] +pub struct Ray { + pos: [f32; 3], + dist: f32, + dir: [f32; 3], + hit: u32, +} + pub struct VoxelRenderer { render_texture: crawl::Texture, render_pipeline: wgpu::RenderPipeline, + raygen_bind_group: wgpu::BindGroup, + raygen_pipeline: wgpu::ComputePipeline, } impl VoxelRenderer { - pub fn new(context: &crawl::Context) -> Result { - let shader = context + pub fn new(context: &crawl::Context, camera: &CameraController) -> Result { + let render_shader = context .device .create_shader_module(wgpu::include_wgsl!(concat!( env!("CARGO_MANIFEST_DIR"), @@ -26,20 +39,68 @@ impl VoxelRenderer { .with_shader_visibility(wgpu::ShaderStages::FRAGMENT | wgpu::ShaderStages::COMPUTE) .build(context)?; - let render_pipeline = crawl::RenderPipelineBuilder::new("Raycast Quad", &shader) + let render_pipeline = crawl::RenderPipelineBuilder::new("Raycast Quad", &render_shader) .with_layout("Draw", &[&render_texture.bind_group_layout], &[]) .with_fragment_targets(&[Some(context.surface_config.format.into())]) .build(context); + let raygen_shader = context + .device + .create_shader_module(wgpu::include_wgsl!(concat!( + env!("CARGO_MANIFEST_DIR"), + "/res/common/raygen.wgsl" + ))); + + let ray_data = vec![Ray::default(); (context.size.width * context.size.height) as usize]; + let mut raygen_buffers = crawl::BulkBufferBuilder::new() + .set_usage(wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST) + .with_init_buffer_bm("Rays", &ray_data) + .build(context); + let raygen_layout = crawl::BindGroupLayoutBuilder::new() + .with_label("Raygen BGL") + .with_entry( + wgpu::ShaderStages::COMPUTE, + wgpu::BindingType::StorageTexture { + access: wgpu::StorageTextureAccess::WriteOnly, + format: render_texture.attributes.format, + view_dimension: wgpu::TextureViewDimension::D2, + }, + None, + ) + .with_uniform_entry(wgpu::ShaderStages::COMPUTE) + .with_rw_storage_entry(wgpu::ShaderStages::COMPUTE) + .build(context); + + let raygen_bind_group = crawl::BindGroupBuilder::new() + .with_label("Raygen BG") + .with_layout(&raygen_layout) + .with_entry(wgpu::BindingResource::TextureView(&render_texture.view)) + .with_entry(camera.get_buffer().as_entire_binding()) + .with_entry(raygen_buffers.remove(0).as_entire_binding()) + .build(context)?; + + let raygen_pipeline = crawl::ComputePipelineBuilder::new("Raygen P", &raygen_shader) + .with_layout("Raygen PL", &[&raygen_layout], &[]) + .build(context); + Ok(Self { render_texture, render_pipeline, + raygen_pipeline, + raygen_bind_group, }) } } impl crawl::Pass for VoxelRenderer { fn execute(&self, encoder: &mut wgpu::CommandEncoder, view: &wgpu::TextureView) { + let size = self.render_texture.attributes.size; + let mut compute_pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor::default()); + compute_pass.set_pipeline(&self.raygen_pipeline); + compute_pass.set_bind_group(0, &self.raygen_bind_group, &[]); + compute_pass.dispatch_workgroups(size.width / 8, size.height / 8, 1); + drop(compute_pass); + let mut render_pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { label: Some("Render Pass"), color_attachments: &[Some(wgpu::RenderPassColorAttachment { @@ -52,5 +113,6 @@ impl crawl::Pass for VoxelRenderer { render_pass.set_pipeline(&self.render_pipeline); render_pass.set_bind_group(0, &self.render_texture.bind_group, &[]); render_pass.draw(0..6, 0..1); + drop(render_pass); } }