render: create light cull compute shader, bind buffers, etc.

lyra-game/src/render/camera.rs

@@ -1,5 +1,3 @@
-use std::{mem, num::NonZeroU64};
-
 use winit::dpi::PhysicalSize;
 
 use crate::{math::{Angle, OPENGL_TO_WGPU_MATRIX}, scene::CameraComponent};

lyra-game/src/render/light/mod.rs

@@ -103,18 +103,21 @@ pub(crate) struct LightUniformBuffers {
     pub buffer: wgpu::Buffer,
     pub bindgroup_layout: wgpu::BindGroupLayout,
     pub bindgroup: wgpu::BindGroup,
-    pub lights_uniform: LightsUniform,
+    pub light_indexes: HashMap<Entity, u32>,
-    pub point_lights: LightBuffer<PointLightUniform>,
+    pub current_light_idx: u32,
-    pub spot_lights: LightBuffer<SpotLightUniform>,
 }
 
 impl LightUniformBuffers {
     pub fn new(device: &wgpu::Device) -> Self {
+        let limits = device.limits();
+        // TODO: check this limit somehow
+        let max_buffer_sizes = (limits.max_uniform_buffer_binding_size as u64) / 2;
+
         let buffer = device.create_buffer(
             &wgpu::BufferDescriptor {
                 label: Some("UBO_Lights"),
-                usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
+                usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
-                size: mem::size_of::<LightsUniform>() as u64,
+                size: max_buffer_sizes,
                 mapped_at_creation: false,
             }
         );
@@ -123,14 +126,16 @@ impl LightUniformBuffers {
             entries: &[
                 wgpu::BindGroupLayoutEntry {
                     binding: 0,
-                    visibility: wgpu::ShaderStages::VERTEX | wgpu::ShaderStages::FRAGMENT,
+                    visibility: wgpu::ShaderStages::FRAGMENT | wgpu::ShaderStages::COMPUTE,
                     ty: wgpu::BindingType::Buffer {
-                        ty: wgpu::BufferBindingType::Uniform,
+                        ty: wgpu::BufferBindingType::Storage {
+                            read_only: true
+                        },
                         has_dynamic_offset: false,
                         min_binding_size: None,
                     },
                     count: None,
-                }
+                },
             ],
             label: Some("BGL_Lights"),
         });
@@ -147,63 +152,167 @@ impl LightUniformBuffers {
                             size: None, // use the full buffer
                         }
                     )
-                }
+                },
             ],
             label: Some("BG_Lights"),
         });
 
-        let point_lights = LightBuffer::new(MAX_LIGHT_COUNT);
-        let spot_lights = LightBuffer::new(MAX_LIGHT_COUNT);
-
         Self {
             buffer,
             bindgroup_layout,
             bindgroup,
-            lights_uniform: LightsUniform::default(),
+            light_indexes: Default::default(),
-            point_lights,
+            current_light_idx: 0,
-            spot_lights,
         }
     }
 
     pub fn update_lights(&mut self, queue: &wgpu::Queue, world_tick: Tick, world: &World) {
+        let mut lights = LightsUniform::default();
+        
         for (entity, point_light, transform, light_epoch, transform_epoch)
                 in world.view_iter::<(Entities, &PointLight, &Transform, TickOf<PointLight>, TickOf<Transform>)>() {
 
-            if !self.point_lights.has_light(entity) || light_epoch == world_tick || transform_epoch == world_tick {
+            // TODO: dont update light every frame
-                let uniform = PointLightUniform::from_bundle(&point_light, &transform);
+            let idx = *self.light_indexes.entry(entity) 
-                self.point_lights.update_or_add(&mut self.lights_uniform.point_lights, entity, uniform);
+                .or_insert_with(|| {
-                //debug!("Updated point light");
+                    let t = self.current_light_idx;
-            }
+                    self.current_light_idx += 1;
+                    t
+                }) as usize;
+
+            let uniform = LightUniform::from_point_light_bundle(&point_light, &transform);
+            lights.data[idx] = uniform;
         }
 
-        for (entity, spot_light, transform, light_epoch, transform_epoch)
+        if let Some((entity, dir_light, transform)) = 
-                in world.view_iter::<(Entities, &SpotLight, &Transform, TickOf<SpotLight>, TickOf<Transform>)>() {
+            world.view_iter::<(Entities, &DirectionalLight, &Transform)>().next() {
-
-            if !self.spot_lights.has_light(entity) || light_epoch == world_tick || transform_epoch == world_tick {
-                let uniform = SpotLightUniform::from_bundle(&spot_light, &transform);
-                self.spot_lights.update_or_add(&mut self.lights_uniform.spot_lights, entity, uniform);
-                //debug!("Updated spot light");
-            }
-        }
-
-        if let Some((dir_light, transform)) = 
-                world.view_iter::<(&DirectionalLight, &Transform)>().next() {
             
-            let uniform = DirectionalLightUniform::from_bundle(&dir_light, &transform);
+            let idx = *self.light_indexes.entry(entity) 
-            self.lights_uniform.directional_light = uniform;
+                .or_insert_with(|| {
+                    let t = self.current_light_idx;
+                    self.current_light_idx += 1;
+                    t
+                }) as usize;
+
+            let uniform = LightUniform::from_directional_bundle(&dir_light, &transform);
+            lights.data[idx] = uniform;
         }
 
-        self.lights_uniform.point_light_count = self.point_lights.buffer_count as u32;
+        lights.light_count = self.light_indexes.len() as u32;
-        self.lights_uniform.spot_light_count = self.spot_lights.buffer_count as u32;
+
-        queue.write_buffer(&self.buffer, 0, bytemuck::cast_slice(&[self.lights_uniform]));
+        // update the light count in the struct
+        queue.write_buffer(&self.buffer, 0, bytemuck::cast_slice(&[lights]));
     }
 }
 
+#[repr(C)]
+#[derive(Default, Debug, Copy, Clone)]
+pub(crate) enum LightType {
+    #[default]
+    Directional = 0,
+    Point = 1,
+    Spotlight = 2,
+}
+
+#[repr(C)]
+#[derive(Default, Debug, Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)]
+pub(crate) struct LightUniform {
+    pub light_type: u32, // LightType
+    pub enabled: u32, // bool
+    pub _padding: [u32; 2],
+    pub position: glam::Vec3,
+    pub _padding2: u32,
+    pub direction: glam::Vec3,
+    pub _padding3: u32,
+    pub color: glam::Vec3,
+    // no padding is needed here since cutoff acts as the padding
+    // that would usually be needed for the vec3
+
+    pub range: f32,
+    pub intensity: f32,
+
+    pub spot_cutoff: f32,
+    pub spot_outer_cutoff: f32,
+
+    pub _padding4: u32,
+}
+
+impl LightUniform {
+    pub fn from_point_light_bundle(light: &PointLight, transform: &Transform) -> Self {
+        Self {
+            light_type: LightType::Point as u32,
+            enabled: true as u32, // TODO
+            _padding: [0; 2],
+            position: transform.translation,
+            _padding2: 0,
+            direction: transform.forward(),
+            _padding3: 0,
+            color: light.color,
+
+            range: 2.0,
+            intensity: 1.0,
+
+            spot_cutoff: 0.0,
+            spot_outer_cutoff: 0.0,
+
+            _padding4: 0,
+        }
+    }
+
+    pub fn from_directional_bundle(light: &DirectionalLight, transform: &Transform) -> Self {
+        Self {
+            light_type: LightType::Directional as u32,
+            enabled: true as u32, // TODO: take from component
+            _padding: [0; 2],
+            position: transform.translation,
+            _padding2: 0,
+            direction: transform.forward(),
+            _padding3: 0,
+            color: light.color,
+
+            range: 0.0,
+            intensity: 0.0,
+
+            spot_cutoff: 0.0,
+            spot_outer_cutoff: 0.0,
+
+            _padding4: 0,
+        }
+    }
+
+    // Create the SpotLightUniform from an ECS bundle
+    /* pub fn from_bundle(light: &SpotLight, transform: &Transform) -> Self {
+        Self {
+            position: transform.translation,
+            _padding: 0,
+            direction: transform.forward(),
+            _padding2: 0,
+            color: light.color,
+            cutoff: light.cutoff.to_radians().cos(),
+            outer_cutoff: light.outer_cutoff.to_radians().cos(),
+            constant: light.constant,
+            linear: light.linear,
+            quadratic: light.quadratic,
+            ambient: light.ambient,
+            diffuse: light.diffuse,
+            specular: light.specular,
+            _padding3: 0,
+        }
+    } */
+}
+
 #[repr(C)]
 #[derive(Default, Debug, Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)]
 pub struct LightsUniform {
-    point_lights: [PointLightUniform; MAX_LIGHT_COUNT],
+    light_count: u32,
-    point_light_count: u32,
+    _padding: [u32; 3],
+    data: [LightUniform; 10], // TODO: make this a dynamic length
+}
+
+/* #[repr(C)]
+#[derive(Default, Debug, Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)]
+pub struct LightsUniform {
+    light_count: u32,
     _padding: [u32; 3],
     spot_lights: [SpotLightUniform; MAX_LIGHT_COUNT],
     spot_light_count: u32,
@@ -352,4 +461,4 @@ impl SpotLightUniform {
             _padding3: 0,
         }
     }
-}
+} */

lyra-game/src/render/light_cull_compute.rs

@@ -0,0 +1,223 @@
+use std::{borrow::Cow, mem, num::NonZeroU32, ptr::NonNull, rc::Rc};
+
+use glam::UVec2;
+use wgpu::{util::DeviceExt, ComputePipeline};
+use winit::dpi::PhysicalSize;
+
+use super::{light::LightUniformBuffers, render_buffer::BufferWrapper, renderer::RenderPass, texture::RenderTexture};
+
+struct LightIndicesGridBuffer {
+    indices_buffer: wgpu::Buffer,
+    grid_texture: wgpu::Texture,
+    grid_texture_view: wgpu::TextureView,
+    bind_group_layout: wgpu::BindGroupLayout,
+    bind_group: wgpu::BindGroup,
+}
+
+pub(crate) struct LightCullCompute {
+    device: Rc<wgpu::Device>,
+    queue: Rc<wgpu::Queue>,
+    pipeline: ComputePipeline,
+    lights: NonNull<LightUniformBuffers>,
+    camera: NonNull<BufferWrapper>,
+    light_indices_grid: LightIndicesGridBuffer,
+    screen_size_buffer: BufferWrapper,
+    depth_tex: NonNull<RenderTexture>,
+}
+
+impl LightCullCompute {
+    fn create_grid(device: &wgpu::Device, screen_size: PhysicalSize<u32>) -> LightIndicesGridBuffer {
+        let limits = device.limits();
+        //let max_buffer_sizes = (limits.max_uniform_buffer_binding_size as u64) / 2;
+
+        /* let light_indices_buffer = device.create_buffer(
+            &wgpu::BufferDescriptor {
+                label: Some("B_LightIndices"),
+                usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
+                size: (mem::size_of::<u32>() * 16 * 16) as u64,
+                mapped_at_creation: false,
+            }
+        ); */
+
+        let light_indices_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
+            label: Some("B_LightIndices"),
+            contents: &[0; mem::size_of::<u32>() * 16 * 16],
+            usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
+        });
+
+        let light_indices_bg_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+            entries: &[
+                wgpu::BindGroupLayoutEntry {
+                    binding: 0,
+                    visibility: wgpu::ShaderStages::COMPUTE | wgpu::ShaderStages::FRAGMENT,
+                    ty: wgpu::BindingType::Buffer {
+                        ty: wgpu::BufferBindingType::Storage {
+                            read_only: false
+                        },
+                        has_dynamic_offset: false,
+                        min_binding_size: None,
+                    },
+                    count: None,
+                },
+                wgpu::BindGroupLayoutEntry {
+                    binding: 1,
+                    visibility: wgpu::ShaderStages::COMPUTE | wgpu::ShaderStages::FRAGMENT,
+                    ty: wgpu::BindingType::StorageTexture {
+                        access: wgpu::StorageTextureAccess::ReadWrite,
+                        format: wgpu::TextureFormat::Rg32Uint, // vec2<uint>
+                        view_dimension: wgpu::TextureViewDimension::D2
+                    },
+                    count: None,
+                }
+            ],
+            label: Some("BGL_LightIndicesGrid"),
+        });
+
+        // TODO: shrink the texture to match the amount of grid cells that the shader actually uses
+        let size = wgpu::Extent3d {
+            width: screen_size.width,
+            height: screen_size.height,
+            depth_or_array_layers: 1,
+        };
+        let grid_texture = device.create_texture(
+            &wgpu::TextureDescriptor {
+                label: Some("Tex_LightGrid"),
+                size,
+                mip_level_count: 1,
+                sample_count: 1,
+                dimension: wgpu::TextureDimension::D2,
+                format: wgpu::TextureFormat::Rg32Uint, // vec2<uint>
+                usage: wgpu::TextureUsages::STORAGE_BINDING,
+                view_formats: &[],
+            }
+        );
+
+        let grid_texture_view = grid_texture.create_view(&wgpu::TextureViewDescriptor {
+            label: Some("TexV_LightGrid"),
+            format: Some(wgpu::TextureFormat::Rg32Uint), // vec2<uint>
+            dimension: Some(wgpu::TextureViewDimension::D2),
+            aspect: wgpu::TextureAspect::All,
+            base_mip_level: 0,
+            mip_level_count: None,
+            base_array_layer: 0,
+            array_layer_count: None,
+        });
+
+        let light_indices_bg = device.create_bind_group(&wgpu::BindGroupDescriptor {
+            layout: &light_indices_bg_layout,
+            entries: &[
+                wgpu::BindGroupEntry {
+                    binding: 0,
+                    resource: wgpu::BindingResource::Buffer(
+                            wgpu::BufferBinding {
+                            buffer: &light_indices_buffer,
+                            offset: 0,
+                            size: None, // the entire light buffer is needed
+                        }
+                    )
+                },
+                wgpu::BindGroupEntry {
+                    binding: 1,
+                    resource: wgpu::BindingResource::TextureView(&grid_texture_view)
+                }
+            ],
+            label: Some("BG_LightIndicesGrid"),
+        });
+
+        LightIndicesGridBuffer {
+            indices_buffer: light_indices_buffer,
+            grid_texture,
+            grid_texture_view,
+            bind_group_layout: light_indices_bg_layout,
+            bind_group: light_indices_bg,
+        }
+    }
+
+    pub fn new(device: Rc<wgpu::Device>, queue: Rc<wgpu::Queue>, screen_size: PhysicalSize<u32>, lights: &LightUniformBuffers, camera_buffers: &BufferWrapper, depth_texture: &mut RenderTexture) -> Self {
+        let screen_size_buffer = BufferWrapper::builder()
+            .buffer_usage(wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST)
+            .label_prefix("ScreenSize")
+            .visibility(wgpu::ShaderStages::COMPUTE)
+            .buffer_dynamic_offset(false)
+            .contents(&[UVec2::new(screen_size.width, screen_size.height)])
+            .finish(&device);
+
+        let shader_src = include_str!("shaders/light_cull.comp.wgsl");
+        let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
+            label: Some("LightCullCompute"),
+            source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(shader_src)),
+        });
+
+        let light_grid = Self::create_grid(&device, screen_size);
+
+        let depth_tex_pair = depth_texture.create_bind_group(&device);
+
+        let layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+            label: Some("PipeLay_LightCull"),
+            bind_group_layouts: &[
+                //&depth_texture.bindgroup_pair.as_ref().unwrap().layout,
+                &depth_tex_pair.layout,
+                camera_buffers.bindgroup_layout().unwrap(),
+                &lights.bindgroup_layout,
+                &light_grid.bind_group_layout,
+                screen_size_buffer.bindgroup_layout().unwrap(),
+            ],
+            push_constant_ranges: &[],
+        });
+
+        let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+            label: Some("Pipe_LightCull"),
+            layout: Some(&layout),
+            module: &shader,
+            entry_point: "cs_main",
+        });
+
+        Self {
+            device,
+            queue,
+            pipeline,
+            lights: NonNull::from(lights),
+            camera: NonNull::from(camera_buffers),
+            light_indices_grid: light_grid,
+            screen_size_buffer,
+            depth_tex: NonNull::from(depth_texture),
+        }
+    }
+
+    pub fn update_screen_size(&self, size: PhysicalSize<u32>) {
+        self.screen_size_buffer.write_buffer(&self.queue, 0,
+            &[UVec2::new(size.width, size.height)]);
+    }
+
+    pub fn compute(&mut self) {
+        //self.queue.write_buffer(&self.light_indices_grid.indices_buffer, 0, &[0; mem::size_of::<u32>() * 16 * 16]);
+
+        let mut encoder = self.device.create_command_encoder(&wgpu::CommandEncoderDescriptor {
+            label: Some("LightCullCompute"),
+        });
+
+        {
+            let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
+                label: Some("Pass_LightCull"),
+            });
+
+            pass.set_pipeline(&self.pipeline);
+            
+            let depth = unsafe { self.depth_tex.as_ref() };
+            pass.set_bind_group(0, depth.bind_group(), &[]);
+
+            let cam = unsafe { self.camera.as_ref() };
+            pass.set_bind_group(1, cam.bindgroup(), &[]);
+
+            let lights = unsafe { self.lights.as_ref() };
+            pass.set_bind_group(2, &lights.bindgroup, &[]);
+
+            pass.set_bind_group(3, &self.light_indices_grid.bind_group, &[]);
+            pass.set_bind_group(4, self.screen_size_buffer.bindgroup(), &[]);
+            
+            pass.dispatch_workgroups(16, 16, 1);
+        }
+        self.queue.submit(std::iter::once(encoder.finish()));
+        self.device.poll(wgpu::Maintain::Wait);
+    }
+}

lyra-game/src/render/mod.rs

@@ -11,4 +11,5 @@ pub mod material;
 pub mod camera;
 pub mod window;
 pub mod transform_buffer_storage;
-pub mod light;
+pub mod light;
+pub mod light_cull_compute;

lyra-game/src/render/render_buffer.rs

@@ -39,12 +39,19 @@ impl BindGroupPair {
             layout,
         }
     }
+
+    pub fn new(bindgroup: wgpu::BindGroup, layout: wgpu::BindGroupLayout) -> Self {
+        Self {
+            bindgroup,
+            layout: Arc::new(layout),
+        }
+    }
 }
 
 pub struct BufferWrapper {
     pub bindgroup_pair: Option<BindGroupPair>,
     pub inner_buf: wgpu::Buffer,
-    pub len: usize,
+    pub len: Option<usize>,
 }
 
 impl BufferWrapper {
@@ -54,7 +61,7 @@ impl BufferWrapper {
         Self {
             bindgroup_pair: bind_group,
             inner_buf: buffer,
-            len: 0,
+            len: Some(0),
         }
     }
 
@@ -64,7 +71,15 @@ impl BufferWrapper {
         Self {
             bindgroup_pair: bind_group,
             inner_buf: buffer,
-            len: 0,
+            len: Some(0),
+        }
+    }
+
+    pub fn from_parts(bind_group: wgpu::BindGroup, bind_group_layout: wgpu::BindGroupLayout, buffer: wgpu::Buffer) -> Self {
+        Self {
+            bindgroup_pair: Some(BindGroupPair::new(bind_group, bind_group_layout)),
+            inner_buf: buffer,
+            len: None,
         }
     }
 
@@ -101,7 +116,7 @@ impl BufferWrapper {
     /// match the layout of this bind group.
     /// 
     /// See [`wgpu::RenderPass::set_bind_group`](https://docs.rs/wgpu/latest/wgpu/struct.RenderPass.html#method.set_bind_group).
-    pub fn bind_at_bind_group<'a, 'b>(
+    pub fn render_pass_bind_at<'a, 'b>(
         &'a self,
         pass: &'b mut wgpu::RenderPass<'a>,
         index: u32,
@@ -112,6 +127,13 @@ impl BufferWrapper {
         );
         pass.set_bind_group(index, &pair.bindgroup, offsets);
     }
+
+    /// Returns the bindgroup of this buffer, panics if the buffer does not have a bindgroup.
+    pub fn bindgroup(&self) -> &wgpu::BindGroup {
+        &self.bindgroup_pair.as_ref().expect(
+            "BufferWrapper is missing bindgroup pair! Cannot set bind group on RenderPass!",
+        ).bindgroup
+    }
 }
 
 /// Struct used for building a BufferWrapper
@@ -274,7 +296,7 @@ impl BufferWrapperBuilder {
         BufferWrapper {
             bindgroup_pair: Some(bg_pair),
             inner_buf: buffer,
-            len: self.count.unwrap_or_default() as usize,
+            len: Some(self.count.unwrap_or_default() as usize),
         }
     }
 }

lyra-game/src/render/renderer.rs

@@ -1,4 +1,5 @@
 use std::collections::{HashMap, VecDeque, HashSet};
+use std::rc::Rc;
 use std::sync::Arc;
 use std::borrow::Cow;
 
@@ -24,6 +25,7 @@ use crate::scene::CameraComponent;
 use super::camera::{RenderCamera, CameraUniform};
 use super::desc_buf_lay::DescVertexBufferLayout;
 use super::light::LightUniformBuffers;
+use super::light_cull_compute::LightCullCompute;
 use super::material::Material;
 use super::render_buffer::BufferWrapper;
 use super::texture::RenderTexture;
@@ -45,6 +47,12 @@ pub trait Renderer {
     fn add_render_pipeline(&mut self, shader_id: u64, pipeline: Arc<FullRenderPipeline>);
 }
 
+pub trait RenderPass {
+    fn prepare(&mut self, main_world: &mut World);
+    fn render(&mut self, encoder: &mut wgpu::CommandEncoder) -> Result<(), wgpu::SurfaceError>;
+    fn on_resize(&mut self, new_size: winit::dpi::PhysicalSize<u32>);
+}
+
 struct MeshBufferStorage {
     buffer_vertex: BufferStorage,
     buffer_indices: Option<(wgpu::IndexFormat, BufferStorage)>,
@@ -68,8 +76,8 @@ pub struct CachedTransform {
 
 pub struct BasicRenderer {
     pub surface: wgpu::Surface,
-    pub device: wgpu::Device,
+    pub device: Rc<wgpu::Device>, // device does not need to be mutable, no need for refcell
-    pub queue: wgpu::Queue,
+    pub queue: Rc<wgpu::Queue>,
     pub config: wgpu::SurfaceConfiguration,
     pub size: winit::dpi::PhysicalSize<u32>,
     pub window: Arc<Window>,
@@ -98,6 +106,8 @@ pub struct BasicRenderer {
     material_buffer: BufferWrapper,
 
     light_buffers: LightUniformBuffers,
+
+    light_cull_compute: LightCullCompute,
 }
 
 impl BasicRenderer {
@@ -123,7 +133,7 @@ impl BasicRenderer {
 
         let (device, queue) = adapter.request_device(
             &wgpu::DeviceDescriptor {
-                features: wgpu::Features::empty(),
+                features: wgpu::Features::TEXTURE_ADAPTER_SPECIFIC_FORMAT_FEATURES,
                 // WebGL does not support all wgpu features.
                 // Not sure if the engine will ever completely support WASM,
                 // but its here just in case
@@ -175,11 +185,11 @@ impl BasicRenderer {
             .buffer_usage(wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST)
             .contents(&[CameraUniform::default()])
             .label_prefix("Camera")
-            .visibility(wgpu::ShaderStages::VERTEX | wgpu::ShaderStages::FRAGMENT)
+            .visibility(wgpu::ShaderStages::VERTEX | wgpu::ShaderStages::FRAGMENT | wgpu::ShaderStages::COMPUTE)
             .buffer_dynamic_offset(false)
             .finish(&device);
 
-        let depth_texture = RenderTexture::create_depth_texture(&device, &config, "Depth Buffer");
+        let mut depth_texture = RenderTexture::create_depth_texture(&device, &config, "Tex_Depth");
 
         // load the default texture
         let bytes = include_bytes!("default_texture.png");
@@ -193,6 +203,10 @@ impl BasicRenderer {
             .contents(&[MaterialUniform::default()])
             .finish(&device);
 
+        let device = Rc::new(device);
+        let queue = Rc::new(queue);
+        let light_cull_compute = LightCullCompute::new(device.clone(), queue.clone(), size, &light_uniform_buffers, &camera_buffer, &mut depth_texture);
+
         let mut s = Self {
             window,
             surface,
@@ -224,6 +238,7 @@ impl BasicRenderer {
 
             light_buffers: light_uniform_buffers,
             material_buffer: mat_buffer,
+            light_cull_compute,
         };
 
         // create the default pipelines
@@ -494,6 +509,8 @@ impl Renderer for BasicRenderer {
         let output = self.surface.get_current_texture()?;
         let view = output.texture.create_view(&wgpu::TextureViewDescriptor::default());
 
+        //self.light_cull_compute.compute();
+
         let mut encoder = self.device.create_command_encoder(&wgpu::CommandEncoderDescriptor {
             label: Some("Basic Renderer's Encoder")
         });
@@ -553,7 +570,7 @@ impl Renderer for BasicRenderer {
                     let offset = TransformBuffers::index_offset(&self.render_limits, transform_indices) as u32;
                     render_pass.set_bind_group(1, bindgroup, &[ offset, offset, ]);
 
-                    self.camera_buffer.bind_at_bind_group(&mut render_pass, 2, &[]);
+                    render_pass.set_bind_group(2, &self.camera_buffer.bindgroup(), &[]);
                     render_pass.set_bind_group(3, &self.light_buffers.bindgroup, &[]);
                     render_pass.set_bind_group(4, &self.material_buffer.bindgroup_pair.as_ref().unwrap().bindgroup, &[]);
         
@@ -590,6 +607,7 @@ impl Renderer for BasicRenderer {
             self.surface.configure(&self.device, &self.config);
             self.depth_buffer_texture = RenderTexture::create_depth_texture(&self.device, &self.config, "Depth Buffer Texture");
             self.inuse_camera.update_aspect_ratio(self.size);
+            self.light_cull_compute.update_screen_size(new_size);
         }
     }
 

lyra-game/src/render/shaders/base.wgsl

@@ -2,6 +2,10 @@
 
 const max_light_count: u32 = 16u;
 
+const light_ty_directional = 0u;
+const light_ty_point = 1u;
+const light_ty_spot = 2u;
+
 struct VertexInput {
     @location(0) position: vec3<f32>,
     @location(1) tex_coords: vec2<f32>,
@@ -21,53 +25,25 @@ struct CameraUniform {
     view_pos: vec3<f32>,
 };
 
-struct PointLight {
+struct Light {
-    position: vec4<f32>,
+    light_ty: u32,
-    color: vec4<f32>,
+    enabled: u32,
-    
-    intensity: f32,
-    constant: f32,
-    linear: f32,
-    quadratic: f32,
-    
-    ambient: f32,
-    diffuse: f32,
-    specular: f32,
-};
 
-struct DirectionalLight {
-    direction: vec3<f32>,
-    color: vec3<f32>,
-    
-    ambient: f32,
-    diffuse: f32,
-    specular: f32,
-};
-
-struct SpotLight {
     position: vec3<f32>,
     direction: vec3<f32>,
     color: vec3<f32>,
     
-    cutoff: f32,
+    range: f32,
-    outer_cutoff: f32,
+    intensity: f32,
 
-    constant: f32,
+    spot_cutoff: f32,
-    linear: f32,
+    spot_outer_cutoff: f32,
-    quadratic: f32,
-
-    ambient: f32,
-    diffuse: f32,
-    specular: f32,
 };
 
 struct Lights {
-    point_lights: array<PointLight, max_light_count>,
+    light_count: u32,
-    point_light_count: u32,
+    data: array<Light>,
-    spot_lights: array<SpotLight, max_light_count>,
+};
-    spot_light_count: u32,
-    directional_light: DirectionalLight,
-}
 
 @group(1) @binding(0)
 var<uniform> u_model_transform: mat4x4<f32>;
@@ -78,7 +54,7 @@ var<uniform> u_model_normal_matrix: mat4x4<f32>;
 var<uniform> u_camera: CameraUniform;
 
 @group(3) @binding(0)
-var<uniform> u_lights: Lights;
+var<storage> u_lights: Lights;
 
 @vertex
 fn vs_main(
@@ -129,7 +105,21 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
     // this needs to be 0.0 for the math
     //u_lights.directional_light.direction.w = 0.0;
 
-    var light_res = blinn_phong_dir_light(in.world_position, in.world_normal, u_lights.directional_light, u_material, specular_color);
+    var light_res = vec3<f32>(0.0);
+
+    for (var i = 0u; i < u_lights.light_count; i++) {
+        var light = u_lights.data[i];
+
+        if (light.light_ty == light_ty_directional) {
+            light_res += blinn_phong_dir_light(in.world_position, in.world_normal, light, u_material, specular_color);
+        } else if (light.light_ty == light_ty_point) {
+            light_res += blinn_phong_point_light(in.world_position, in.world_normal, light, u_material, specular_color);
+        } else if (light.light_ty == light_ty_spot) {
+            light_res += blinn_phong_spot_light(in.world_position, in.world_normal, light, u_material, specular_color);
+        }
+    }
+
+    /*var light_res = blinn_phong_dir_light(in.world_position, in.world_normal, u_lights.directional_light, u_material, specular_color);
 
     for (var i = 0u; i < u_lights.point_light_count; i++) {
         light_res += blinn_phong_point_light(in.world_position, in.world_normal, u_lights.point_lights[i], u_material, specular_color);
@@ -137,14 +127,14 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
 
     for (var i = 0u; i < u_lights.spot_light_count; i++) {
         light_res += blinn_phong_spot_light(in.world_position, in.world_normal, u_lights.spot_lights[i], u_material, specular_color);
-    }
+    }*/
 
     let light_object_res = light_res * (object_color.xyz/* * u_material.diffuse.xyz*/);
 
     return vec4<f32>(light_object_res, object_color.a);
 }
 
-fn blinn_phong_dir_light(world_pos: vec3<f32>, world_norm: vec3<f32>, dir_light: DirectionalLight, material: Material, specular_factor: vec3<f32>) -> vec3<f32> {
+fn blinn_phong_dir_light(world_pos: vec3<f32>, world_norm: vec3<f32>, dir_light: Light, material: Material, specular_factor: vec3<f32>) -> vec3<f32> {
     let light_color = dir_light.color.xyz;
     let camera_view_pos = u_camera.view_pos.xyz;
 
@@ -166,14 +156,14 @@ fn blinn_phong_dir_light(world_pos: vec3<f32>, world_norm: vec3<f32>, dir_light:
     var specular_color = specular_strength * (light_color * specular_factor);
     //// end of specular ////
     
-    ambient_color *= dir_light.ambient;
+    /*ambient_color *= dir_light.ambient;
     diffuse_color *= dir_light.diffuse;
-    specular_color *= dir_light.specular;
+    specular_color *= dir_light.specular;*/
 
     return ambient_color + diffuse_color + specular_color;
 }
 
-fn blinn_phong_point_light(world_pos: vec3<f32>, world_norm: vec3<f32>, point_light: PointLight, material: Material, specular_factor: vec3<f32>) -> vec3<f32> {
+fn blinn_phong_point_light(world_pos: vec3<f32>, world_norm: vec3<f32>, point_light: Light, material: Material, specular_factor: vec3<f32>) -> vec3<f32> {
     let light_color = point_light.color.xyz;
     let light_pos = point_light.position.xyz;
     let camera_view_pos = u_camera.view_pos.xyz;
@@ -196,8 +186,14 @@ fn blinn_phong_point_light(world_pos: vec3<f32>, world_norm: vec3<f32>, point_li
     var specular_color = specular_strength * (light_color * specular_factor);
     //// end of specular ////
 
-    //// point light attenuation ////
+    // TODO: Point light range
     let distance = length(light_pos - world_pos);
+    // TODO: make smoothness in this a configurable value
+    // 0.75 is the smoothness or falloff
+    let attenuation = 1.0 - smoothstep(point_light.range * 0.75, point_light.range, distance);
+
+    //// point light attenuation ////
+    /*let distance = length(light_pos - world_pos);
     let attenuation = 1.0 / (point_light.constant + point_light.linear * distance + 
     	point_light.quadratic * (distance * distance));
 
@@ -205,58 +201,15 @@ fn blinn_phong_point_light(world_pos: vec3<f32>, world_norm: vec3<f32>, point_li
     
     ambient_color *= point_light.ambient * attenuation;
     diffuse_color *= point_light.diffuse * attenuation;
-    specular_color *= point_light.specular * attenuation;
+    specular_color *= point_light.specular * attenuation;*/
+
+    ambient_color *= attenuation;
+    diffuse_color *= attenuation;
+    specular_color *= attenuation;
 
     return (ambient_color + diffuse_color + specular_color) * point_light.intensity;
 }
 
-fn blinn_phong_spot_light(world_pos: vec3<f32>, world_norm: vec3<f32>, spot_light: SpotLight, material: Material, specular_factor: vec3<f32>) -> vec3<f32> {
+fn blinn_phong_spot_light(world_pos: vec3<f32>, world_norm: vec3<f32>, spot_light: Light, material: Material, specular_factor: vec3<f32>) -> vec3<f32> {
-    let light_color = spot_light.color;//.xyz;
+    return vec3<f32>(0.0); // TODO
-    let light_pos = spot_light.position.xyz;
-    let camera_view_pos = u_camera.view_pos.xyz;
-
-    let light_dir = normalize(spot_light.position - world_pos);
-    
-    
-    //if (theta > spot_light.cutoff) {
-        var ambient_color = light_color * material.ambient.xyz * material.diffuse.xyz;
-
-        //// diffuse ////
-        //let light_dir = normalize(light_pos - world_pos);
-
-        let diffuse_strength = max(dot(world_norm, light_dir), 0.0);
-        var diffuse_color = light_color * (diffuse_strength * material.diffuse.xyz);
-        //// end of diffuse ////
-
-        //// specular ////
-        let view_dir = normalize(camera_view_pos - world_pos);
-        let half_dir = normalize(view_dir + light_dir);
-
-        let specular_strength = pow(max(dot(world_norm, half_dir), 0.0), material.shininess);
-        var specular_color = specular_strength * (light_color * specular_factor);
-        //// end of specular ////
-        
-        //// spot light soft edges ////
-        let theta = dot(light_dir, normalize(-spot_light.direction));
-        let epsilon = spot_light.cutoff - spot_light.outer_cutoff;
-        let intensity = clamp((theta - spot_light.outer_cutoff) / epsilon, 0.0, 1.0);
-        //diffuse_color *= intensity;
-        //specular_color *= intensity;
-        //// end of spot light soft edges ////
-
-        //// spot light attenuation ////
-        let distance = length(light_pos - world_pos);
-        let attenuation = 1.0 / (spot_light.constant + spot_light.linear * distance + 
-            spot_light.quadratic * (distance * distance));
-
-        ambient_color *= attenuation * intensity * spot_light.ambient;
-        diffuse_color *= attenuation * intensity * spot_light.diffuse;
-        specular_color *= attenuation * intensity * spot_light.specular;
-        //// end of spot light attenuation ////
-
-
-        return /*ambient_color +*/ diffuse_color + specular_color;
-    /*} else {
-        return vec3<f32>(0.0);
-    }*/
 }

lyra-game/src/render/shaders/light_cull.comp.wgsl

@@ -0,0 +1,248 @@
+const block_size: i32 = 16;
+const max_tile_visible_lights: u32 = 1024u;
+
+const light_ty_directional = 0u;
+const light_ty_point = 1u;
+const light_ty_spot = 2u;
+
+// Possible computer shader inputs:
+//
+// local_invocation_id
+// workgroup_id
+// global_invocation_id
+// num_workgroups
+// local_invocation_index
+
+struct CameraUniform {
+    view_mat: mat4x4<f32>,
+    view_proj: mat4x4<f32>,
+    view_pos: vec3<f32>,
+};
+
+struct Light {
+    light_ty: u32,
+    enabled: u32,
+
+    position: vec3<f32>,
+    direction: vec3<f32>,
+    color: vec3<f32>,
+    
+    range: f32,
+    intensity: f32,
+
+    spot_cutoff: f32,
+    spot_outer_cutoff: f32,
+};
+
+struct Lights {
+    light_count: u32,
+    data: array<Light>,
+};
+
+var<workgroup> wg_min_depth: atomic<u32>;
+var<workgroup> wg_max_depth: atomic<u32>;
+var<workgroup> wg_frustum_planes: array<vec4<f32>, 6>;
+
+// index list of visible light sources for this tile
+var<workgroup> wg_visible_light_indices: array<u32, max_tile_visible_lights>;
+var<workgroup> wg_visible_light_count: atomic<u32>;
+
+//var<workgroup> view_projection: mat4x4;
+
+@group(0) @binding(0)
+var t_depthmap: texture_2d<f32>;
+@group(0) @binding(1)
+var s_depthmap: sampler;
+
+@group(1) @binding(0)
+var<uniform> u_camera: CameraUniform;
+
+@group(2) @binding(0)
+var<storage, read> u_lights: Lights;
+
+@group(3) @binding(0)
+var<storage, read_write> u_light_indices: array<u32>;
+/*@group(3) @binding(1)
+var<uniform> u_light_grid: array<array<vec2<u32>>>;*/
+
+@group(3) @binding(1)
+var t_light_grid: texture_storage_2d<rg32uint, read_write>; // rg32uint = vec2<u32> or vec4<u32>(r, g, 0.0, 1.0) 
+
+@group(4) @binding(0)
+var<uniform> u_screen_size: vec2<u32>;
+
+@compute
+@workgroup_size(16, 16, 1)
+fn cs_main(
+    @builtin(local_invocation_id) local_invocation_id: vec3<u32>,
+    @builtin(workgroup_id) workgroup_id: vec3<u32>,
+    @builtin(global_invocation_id) global_invocation_id: vec3<u32>,
+    @builtin(num_workgroups) num_workgroups: vec3<u32>,
+    @builtin(local_invocation_index) local_invocation_index: u32,
+) {
+    //var location = vec2<i32>(global_invocation_id.xy);
+    var item_id = vec2<i32>(local_invocation_id.xy);
+    var tile_id = vec2<i32>(workgroup_id.xy);
+    var tile_number = vec2<i32>(num_workgroups.xy);
+    var index = tile_id.y * tile_number.x + tile_id.x;
+
+    // Initialize some shared global values for depth and light count
+    if (local_invocation_index == 0u) {
+        wg_min_depth = 0xFFFFFFFu;
+        wg_max_depth = 0u;
+        wg_visible_light_count = 0u;
+    }
+
+    workgroupBarrier();
+
+    // step 1: calculate the minimum and maximum depth values for this tile (using the depth map)
+    var tex_coord = vec2<f32>(global_invocation_id.xy);
+    //var depth_float: f32 = textureSample(t_depthmap, s_depthmap, tex_coord).r;
+    var depth_float = 0.0;
+    // bitcast the floating depth to u32 for atomic comparisons between threads
+    var depth_uint: u32 = bitcast<u32>(depth_float);
+
+    // step 2: find the minimum and max depth for this tile.
+    // atomically update the workgroup depth
+    atomicMin(&wg_min_depth, depth_uint);
+    atomicMax(&wg_max_depth, depth_uint);
+
+    // convert them back into floats
+    var min_depth: f32 = bitcast<f32>(wg_min_depth);
+    var max_depth: f32 = bitcast<f32>(wg_max_depth);
+
+    workgroupBarrier();
+    
+    // Create the frustum planes that will be used for this time
+    if (local_invocation_index == 0u) {
+        var negative_step = (2.0 * vec2<f32>(tile_id)) / vec2<f32>(tile_number);
+        var positive_step = (2.0 * vec2<f32>(tile_id) + vec2<f32>(1.0, 1.0)) / vec2<f32>(tile_number);
+
+        // z in the vec4 is the distance from the center of the tile
+        wg_frustum_planes[0] = vec4<f32>(1.0, 0.0, 0.0, 1.0 - negative_step.x); // left
+        wg_frustum_planes[1] = vec4<f32>(-1.0, 0.0, 0.0, -1.0 + positive_step.x); // right
+        wg_frustum_planes[2] = vec4<f32>(0.0, -1.0, 0.0, 1.0 - negative_step.y); // bottom
+        wg_frustum_planes[3] = vec4<f32>(0.0, -1.0, 0.0, -1.0 + positive_step.y); // top
+        wg_frustum_planes[4] = vec4<f32>(0.0, 0.0, -1.0, -min_depth); // near plane
+        wg_frustum_planes[5] = vec4<f32>(0.0, 0.0, 1.0, max_depth); // far plane
+
+        // convert the side and top planes from clip to view space
+        for (var i = 0u; i < 4u; i++) {
+            wg_frustum_planes[i] *= u_camera.view_proj;
+            wg_frustum_planes[i] /= length(wg_frustum_planes[i].xyz);
+        }
+
+        // convert near and far planes from clip to view space
+        wg_frustum_planes[4] *= u_camera.view_mat;
+        wg_frustum_planes[4] /= length(wg_frustum_planes[4].xyz);
+        wg_frustum_planes[5] *= u_camera.view_mat;
+        wg_frustum_planes[5] /= length(wg_frustum_planes[5].xyz);
+    }
+
+    workgroupBarrier();
+
+    // Step 3: cull lights
+    
+    // Process the lights detecting which ones to cull for this tile.
+    // Processes 256 lights simultaniously, each on a thread in the workgroup. Requires multiple
+    // iterations for more lights.
+    var thread_count = u32(block_size * block_size);
+    var pass_count = (u_lights.light_count + thread_count - 1u) / thread_count;
+    for (var i = 0u; i < pass_count; i++) {
+        // find the light index to check on this thread, make sure we're not trying to test
+        // for more lights than we have.
+        var light_index = i * thread_count + local_invocation_index;
+        if (light_index > u_lights.light_count) {
+            break;
+        }
+
+        var light = u_lights.data[light_index];
+        var position = light.position;
+        var radius = light.range;
+
+        if (light.light_ty != light_ty_directional
+                && sphere_inside_frustrum(wg_frustum_planes, position, radius)) {
+            // TODO: add the light to the transparent geometry list
+
+            // TODO: spotlights
+            if (!sphere_inside_plane(position, radius, wg_frustum_planes[4])) {
+                var offset: u32 = wg_visible_light_count;
+
+                if (offset < max_tile_visible_lights) {
+                    atomicAdd(&wg_visible_light_count, 1u);
+                    wg_visible_light_indices[offset] = light_index;
+                }
+            }
+        }
+    }
+
+    workgroupBarrier();
+
+    // Update the global memory with the visible light buffer.
+
+    // first update the light grid on the first thread
+    if (local_invocation_index == 0u) {
+        var offset = u32(index) * max_tile_visible_lights; // index in the global light list
+        //t_light_grid[workgroup_id.x][workgroup_id.y] = vec2<f32>(offset, wg_visible_light_count);
+        textureStore(t_light_grid, workgroup_id.xy, vec4<u32>(offset, wg_visible_light_count, 0u, 1u));
+        
+        // TODO: update transparent light grid
+
+        /*var offset = index * max_tile_visible_lights; // position in the global light buffer
+        // update the light
+        for (var i = 0u; i < wg_visible_light_count; i++) {
+            //u_visible_light_indices[offset + i] = wg_visible_light_indices[i];
+        }
+
+        if (wg_visible_light_count != 1024) {
+            // Mark the end of the visible lights for this tile
+            u_visible_light_indices[offset + wg_visible_light_count] = -1;
+        }*/
+    }
+
+    workgroupBarrier();
+
+    // now update the light index list on all threads.
+    var indices_offset = u32(index) * max_tile_visible_lights;
+    //var pass_count = (wg_visible_light_count + thread_count - 1) / thread_count;
+    for (var i = 0u; i < pass_count; i++) {
+        // find the light index to check on this thread, make sure we're not trying to test
+        // for more lights than we have.
+        //var light_index: u32 = i * thread_count + local_invocation_index;
+        /*if (light_index > u_lights.light_count) {
+            u_visible_light_indices
+            break;
+        }*/
+        
+        var offset = indices_offset + i;
+        if (offset >= wg_visible_light_count) {
+            // stop if we're over the over the amount of lights we saw
+            break;
+        }
+
+        u_light_indices[offset] = wg_visible_light_indices[i];
+    }
+}
+
+fn sphere_inside_frustrum(frustum: array<vec4<f32>, 6>, sphere_origin: vec3<f32>, radius: f32) -> bool {
+    // to be able to index this array with a non-const value,
+    // it must be defined as a var
+    var frustum_v = frustum;
+
+    // only check the sides of the frustum
+    for (var i = 0u; i < 4u; i++) {
+        if (!sphere_inside_plane(sphere_origin, radius, frustum_v[i])) {
+            return false;
+        }
+    }
+
+    return true;
+}
+
+/// Check if the sphere is fully behind (i.e., inside the negative half-space of) a plane.
+///
+/// Source: Real-time collision detection, Christer Ericson (2005)
+///   (https://www.3dgep.com/forward-plus/#light-culling-compute-shader)
+fn sphere_inside_plane(sphere_origin: vec3<f32>, radius: f32, plane: vec4<f32>) -> bool {
+    return dot(plane.xyz, sphere_origin) - plane.w < -radius;
+}

lyra-game/src/render/texture.rs

@@ -297,6 +297,66 @@ impl RenderTexture {
         }
     }
 
+    /// Creates a bind group for this texture and returns a borrow to the [`BindGroupPair`]
+    /// 
+    /// This does not create a new bind group if the texture already has one.
+    /// The view dimension will be the same as the texture dimension.
+    pub fn create_bind_group(&mut self, device: &wgpu::Device) -> &BindGroupPair {
+        if self.bindgroup_pair.is_some() {
+            // could not use an if-let here due to the borrow checker thinking
+            // that there was multiple borrows to self.bindgroup_pair
+            return self.bindgroup_pair.as_ref().unwrap();
+        }
+
+        let view_dim = match self.inner_texture.dimension() {
+            wgpu::TextureDimension::D1 => wgpu::TextureViewDimension::D1,
+            wgpu::TextureDimension::D2 => wgpu::TextureViewDimension::D2,
+            wgpu::TextureDimension::D3 => wgpu::TextureViewDimension::D3,
+        };
+
+        let layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+            entries: &[
+                wgpu::BindGroupLayoutEntry {
+                    binding: 0,
+                    visibility: wgpu::ShaderStages::COMPUTE | wgpu::ShaderStages::FRAGMENT,
+                    ty: wgpu::BindingType::Texture {
+                        sample_type: wgpu::TextureSampleType::Depth,
+                        view_dimension: view_dim,
+                        multisampled: false
+                    },
+                    count: None,
+                },
+                wgpu::BindGroupLayoutEntry {
+                    binding: 1,
+                    visibility: wgpu::ShaderStages::COMPUTE | wgpu::ShaderStages::FRAGMENT,
+                    ty: wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Comparison),
+                    count: None,
+                }
+            ],
+            label: Some("BGL_Texture"),
+        });
+
+        let bg = device.create_bind_group(&wgpu::BindGroupDescriptor {
+            layout: &layout,
+            entries: &[
+                wgpu::BindGroupEntry {
+                    binding: 0,
+                    resource: wgpu::BindingResource::TextureView(&self.view)
+                },
+                wgpu::BindGroupEntry {
+                    binding: 1,
+                    resource: wgpu::BindingResource::Sampler(&self.sampler)
+                }
+            ],
+            label: Some("BG_Texture"),
+        });
+
+        let pair = BindGroupPair::new(bg, layout);
+
+        self.bindgroup_pair = Some(pair);
+        self.bindgroup_pair.as_ref().unwrap()
+    }
+
     /// Returns the bind group stored inside the bind group pair.
     /// 
     /// Panics: