Add raygen shader

res/common/raygen.wgsl

@@ -0,0 +1,43 @@
+@group(0) @binding(0) var output: texture_storage_2d<rgba8unorm, write>;
+@group(0) @binding(1) var<uniform> camera: Camera;
+@group(0) @binding(2) var<storage, read_write> rays: array<Ray>;
+
+struct Camera {
+    projection: mat4x4<f32>,
+    view: mat4x4<f32>,
+    pos: vec3<f32>,
+};
+
+struct Ray {
+    pos: vec3<f32>,
+    dist: f32,
+    dir: vec3<f32>,
+    hit: u32,
+}
+
+@compute @workgroup_size(8, 8, 1)
+fn compute(@builtin(global_invocation_id) global_id: vec3<u32>) {
+    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<f32>(img_coord) / vec2<f32>(img_dims);
+    let screen_pos = img_coord_frac * 2.0 - vec2<f32>(1.0);
+    var ray_eye = camera.projection * vec4<f32>(screen_pos, -1.0, 0.0);
+    ray_eye = vec4<f32>(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>(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);
+}

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<Box<dyn crawl::Pass>> = 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:?}");

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<Self> {
+    pub fn new(context: &crawl::Context, camera: &CameraController) -> Result<Self> {
-        let shader = context
+        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);
     }
 }