render: resize light grid with window, improve light buffer, add spot lights to the light cull compute

Spot lights are buggy. They get culled when they shouldn't be, maybe still an issue with the light grid :(

examples/testbed/src/main.rs

@@ -127,7 +127,7 @@ async fn main() {
             light_tran.rotate_y(math::Angle::Degrees(25.0));
             world.spawn((
                 DirectionalLight {
-                    enabled: true,
+                    enabled: false,
                     color: Vec3::ONE,
                     intensity: 0.35
                     //..Default::default()
@@ -174,6 +174,25 @@ async fn main() {
                 ),
                 cube_mesh.clone(),
             ));
+
+            /* world.spawn((
+                SpotLight {
+                    enabled: true,
+                    color: Vec3::new(1.0, 0.0, 0.0),
+                    intensity: 1.0,
+                    range: 1.5,
+                    ..Default::default()
+                },
+                Transform::new(
+                    Vec3::new(0.0, 0.2, -1.5),
+                    //Vec3::new(-5.0, 1.0, -0.28),
+                    //Vec3::new(-10.0, 0.94, -0.28), 
+
+                    Quat::IDENTITY,
+                    Vec3::new(0.15, 0.15, 0.15),
+                ),
+                cube_mesh.clone(),
+            )); */
         }
 
         /* {

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

@@ -6,7 +6,7 @@ use lyra_ecs::{Entity, Tick, World, query::{Entities, TickOf}};
 pub use point::*;
 pub use spotlight::*;
 
-use std::{collections::{HashMap, VecDeque}, marker::PhantomData};
+use std::{collections::{HashMap, VecDeque}, marker::PhantomData, mem};
 
 use crate::math::Transform;
 
@@ -103,13 +103,14 @@ pub(crate) struct LightUniformBuffers {
     pub buffer: wgpu::Buffer,
     pub bind_group_pair: BindGroupPair,
     pub light_indexes: HashMap<Entity, u32>,
+    dead_indices: VecDeque<u32>,
     pub current_light_idx: u32,
 }
 
 impl LightUniformBuffers {
     pub fn new(device: &wgpu::Device) -> Self {
         let limits = device.limits();
-        // TODO: check this limit somehow
+        // TODO: ensure we dont write over this limit
         let max_buffer_sizes = (limits.max_uniform_buffer_binding_size as u64) / 2;
 
         let buffer = device.create_buffer(
@@ -161,45 +162,88 @@ impl LightUniformBuffers {
             bind_group_pair: BindGroupPair::new(bindgroup, bindgroup_layout),
             light_indexes: Default::default(),
             current_light_idx: 0,
+            dead_indices: VecDeque::new(),
         }
     }
 
+    /// Returns the index for the entity, and if this index is new
+    fn get_index_for(&mut self, missed: &mut HashMap<Entity, u32>, entity: Entity) -> (bool, u32) {
+        let idx = missed.remove(&entity)
+            .map(|v| (false, v))
+            .or_else(||
+                self.dead_indices.pop_front()
+                    .map(|v| (true, v))
+            )
+            .unwrap_or_else(|| {
+                let t = self.current_light_idx;
+                self.current_light_idx += 1;
+                (true, t)
+            });
+        idx
+    }
+
     pub fn update_lights(&mut self, queue: &wgpu::Queue, world_tick: Tick, world: &World) {
-        let mut lights = LightsUniform::default();
+        // used to detect what lights were removed
+        let mut missed_lights: HashMap<Entity, u32> = self.light_indexes.drain().collect();
         
         for (entity, point_light, transform, light_epoch, transform_epoch)
                 in world.view_iter::<(Entities, &PointLight, &Transform, TickOf<PointLight>, TickOf<Transform>)>() {
 
-            // TODO: dont update light every frame
+            let (new, idx) = self.get_index_for(&mut missed_lights, entity);
-            let idx = *self.light_indexes.entry(entity) 
+            self.light_indexes.insert(entity, idx);
-                .or_insert_with(|| {
-                    let t = self.current_light_idx;
-                    self.current_light_idx += 1;
-                    t
-                }) as usize;
 
+            if new || light_epoch == world_tick || transform_epoch == world_tick {
                 let uniform = LightUniform::from_point_light_bundle(&point_light, &transform);
-            lights.data[idx] = uniform;
+                
+                let offset = mem::size_of::<u32>() * 4 + mem::size_of::<LightUniform>() * idx as usize;
+                queue.write_buffer(&self.buffer, offset as _, bytemuck::cast_slice(&[uniform]));
+            }
         }
 
-        if let Some((entity, dir_light, transform)) = 
+        for (entity, spot_light, transform, light_epoch, transform_epoch)
-            world.view_iter::<(Entities, &DirectionalLight, &Transform)>().next() {
+                in world.view_iter::<(Entities, &SpotLight, &Transform, TickOf<SpotLight>, TickOf<Transform>)>() {
 
-            let idx = *self.light_indexes.entry(entity) 
+            let (new, idx) = self.get_index_for(&mut missed_lights, entity);
-                .or_insert_with(|| {
+            self.light_indexes.insert(entity, idx);
-                    let t = self.current_light_idx;
-                    self.current_light_idx += 1;
-                    t
-                }) as usize;
 
+            if new || light_epoch == world_tick || transform_epoch == world_tick {
+                let uniform = LightUniform::from_spot_light_bundle(&spot_light, &transform);
+                
+                let offset = mem::size_of::<u32>() * 4 + mem::size_of::<LightUniform>() * idx as usize;
+                queue.write_buffer(&self.buffer, offset as _, bytemuck::cast_slice(&[uniform]));
+            }
+        }
+
+        for (entity, dir_light, transform, light_epoch, transform_epoch)
+                in world.view_iter::<(Entities, &DirectionalLight, &Transform, TickOf<DirectionalLight>, TickOf<Transform>)>() {
+            
+            let (new, idx) = self.get_index_for(&mut missed_lights, entity);
+            self.light_indexes.insert(entity, idx);
+            
+            if new || light_epoch == world_tick || transform_epoch == world_tick {
                 let uniform = LightUniform::from_directional_bundle(&dir_light, &transform);
-            lights.data[idx] = uniform;
+                let offset = mem::size_of::<u32>() * 4 + mem::size_of::<LightUniform>() * idx as usize;
+                queue.write_buffer(&self.buffer, offset as _, bytemuck::cast_slice(&[uniform]));
+            }
         }
 
-        lights.light_count = self.light_indexes.len() as u32;
+        // anything left in missed_lights were lights that were deleted
+        let len = missed_lights.len();
+        self.dead_indices.reserve(len);
 
-        // update the light count in the struct
+        for (_, v) in missed_lights.drain() {
-        queue.write_buffer(&self.buffer, 0, bytemuck::cast_slice(&[lights]));
+            // write zeros in place of this now dead light, the enabled boolean will be set to false
+            let mut zeros = Vec::new();
+            zeros.resize(mem::size_of::<LightUniform>(), 0u32);
+
+            let offset = mem::size_of::<u32>() * 4 + mem::size_of::<LightUniform>() * v as usize;
+            queue.write_buffer(&self.buffer, offset as _, bytemuck::cast_slice(zeros.as_slice()));
+
+            self.dead_indices.push_back(v);
+        }
+
+        // update the amount of lights, then the array of lights
+        queue.write_buffer(&self.buffer, 0, bytemuck::cast_slice(&[self.current_light_idx as u32]));
     }
 }
 
@@ -227,8 +271,8 @@ pub(crate) struct LightUniform {
     pub intensity: f32,
     pub smoothness: f32,
 
-    pub spot_cutoff: f32,
+    pub spot_cutoff_rad: f32,
-    pub spot_outer_cutoff: f32,
+    pub spot_outer_cutoff_rad: f32,
 }
 
 impl LightUniform {
@@ -244,8 +288,8 @@ impl LightUniform {
             intensity: light.intensity,
             smoothness: light.smoothness,
 
-            spot_cutoff: 0.0,
+            spot_cutoff_rad: 0.0,
-            spot_outer_cutoff: 0.0,
+            spot_outer_cutoff_rad: 0.0,
 
         }
     }
@@ -262,190 +306,27 @@ impl LightUniform {
             intensity: light.intensity,
             smoothness: 0.0,
 
-            spot_cutoff: 0.0,
+            spot_cutoff_rad: 0.0,
-            spot_outer_cutoff: 0.0,
+            spot_outer_cutoff_rad: 0.0,
         }
     }
 
     // Create the SpotLightUniform from an ECS bundle
-    /* pub fn from_bundle(light: &SpotLight, transform: &Transform) -> Self {
+    pub fn from_spot_light_bundle(light: &SpotLight, transform: &Transform) -> Self {
         Self {
+            light_type: LightType::Spotlight as u32,
+            enabled: light.enabled as u32,
             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)]
+            range: light.range,
-#[derive(Default, Debug, Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)]
-pub struct LightsUniform {
-    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,
-    _padding2: [u32; 3],
-    directional_light: DirectionalLightUniform,
-}
-
-#[repr(C)]
-#[derive(Default, Debug, Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)]
-pub struct PointLightUniform {
-    /// The position of the light
-    /// vec4 is used here for gpu padding, w is ignored in the shader
-    pub position: glam::Vec4,
-    
-    /// The color of the light
-    /// vec4 is used here for gpu padding, w is ignored in the shader
-    pub color: glam::Vec4,
-
-    /// The intensity of the light
-    /// This works by just multiplying the result of the lighting
-    /// calculations by this scalar
-    pub intensity: f32,
-
-    /// The constant used in the quadratic attenuation calculation. Its best to leave this at 1.0
-    pub constant: f32,
-
-    /// The linear factor used in the quadratic attenuation calculation.
-    pub linear: f32,
-
-    /// The quadratic factor used in the quadratic attenuation calculation.
-    pub quadratic: f32,
-
-    pub ambient: f32,
-    pub diffuse: f32,
-    pub specular: f32,
-
-    pub _padding: u32,
-}
-
-impl PointLightUniform {
-    /// Create the PointLightUniform from an ECS bundle
-    pub fn from_bundle(light: &PointLight, transform: &Transform) -> Self {
-        Self {
-            position: glam::Vec4::new(transform.translation.x, transform.translation.y, transform.translation.z, 0.0),
-            //_padding: 0,
-            color: glam::Vec4::new(light.color.x, light.color.y, light.color.z, 0.0),
-            //_padding2: 0,
             intensity: light.intensity,
-            constant: light.constant,
+            smoothness: light.smoothness,
-            linear: light.linear,
-            quadratic: light.quadratic,
 
-            ambient: light.ambient,
+            spot_cutoff_rad: light.cutoff.to_radians(),
-            diffuse: light.diffuse,
+            spot_outer_cutoff_rad: light.outer_cutoff.to_radians(),
-            specular: light.specular,
-
-            _padding: 0,
         }
     }
 }
 
-#[repr(C)]
-#[derive(Default, Debug, Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)]
-pub struct DirectionalLightUniform {
-    /// The direction of the light
-    pub direction: glam::Vec3,
-    // gpu padding
-    pub _padding: u32,
-    /// The color of the light
-    pub color: glam::Vec3,
-    // no padding is needed here since ambient acts as the padding
-    // that would usually be needed for the vec3
-
-    /// The scalar of the ambient light created by this caster.
-    pub ambient: f32,
-    /// The scalar of the diffuse light created by this caster.
-    pub diffuse: f32,
-    /// The scalar of the specular reflections created by this caster.
-    pub specular: f32,
-
-    pub _padding2: [u32; 2],
-}
-
-impl DirectionalLightUniform {
-    /// Create the DirectionalLightUniform from an ECS bundle
-    pub fn from_bundle(light: &DirectionalLight, transform: &Transform) -> Self {
-        //transform.forward()
-        Self {
-            direction: transform.forward(),
-            _padding: 0,
-            color: light.color,
-            ambient: light.ambient,
-            diffuse: light.diffuse,
-            specular: light.specular,
-            _padding2: [0; 2],
-        }
-    }
-}
-
-#[repr(C)]
-#[derive(Default, Debug, Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)]
-pub(crate) struct SpotLightUniform {
-    pub position: glam::Vec3,
-    pub _padding: u32,
-    pub direction: glam::Vec3,
-    pub _padding2: 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 cutoff: f32,
-    pub outer_cutoff: f32,
-
-    /// The constant used in the quadratic attenuation calculation. Its best to leave this at 1.0
-    pub constant: f32,
-
-    /// The linear factor used in the quadratic attenuation calculation.
-    pub linear: f32,
-
-    /// The quadratic factor used in the quadratic attenuation calculation.
-    pub quadratic: f32,
-
-    pub ambient: f32,
-    pub diffuse: f32,
-    pub specular: f32,
-    pub _padding3: u32,
-}
-
-impl SpotLightUniform {
-    /// 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,
-        }
-    }
-} */

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

@@ -4,15 +4,26 @@ use crate::math;
 
 #[derive(Debug, Clone, Component)]
 pub struct SpotLight {
+    pub enabled: bool,
     pub color: glam::Vec3,
+    pub range: f32,
+    pub intensity: f32,
+    pub smoothness: f32,
     pub cutoff: math::Angle,
     pub outer_cutoff: math::Angle,
-
+}
-    pub constant: f32,
+
-    pub linear: f32,
+impl Default for SpotLight {
-    pub quadratic: f32,
+    fn default() -> Self {
-
+        Self {
-    pub ambient: f32,
+            enabled: true,
-    pub diffuse: f32,
+            color: glam::Vec3::new(1.0, 1.0, 1.0),
-    pub specular: f32,
+            range: 1.0,
+            intensity: 1.0,
+            smoothness: 0.75,
+
+            cutoff: math::Angle::Degrees(45.0),
+            outer_cutoff: math::Angle::Degrees(45.0),
+        }
+    }
 }

lyra-game/src/render/light_cull_compute.rs

@@ -83,7 +83,6 @@ impl LightCullCompute {
             label: Some("BGL_LightIndicesGrid"),
         });
 
-        // TODO: resize texture when screen is resized
         let size = wgpu::Extent3d {
             width: workgroup_size.x,
             height: workgroup_size.y,
@@ -168,7 +167,8 @@ impl LightCullCompute {
             source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(shader_src)),
         });
 
-        let workgroup_size = glam::UVec2::new((screen_size.width as f32 / 16.0).ceil() as u32, (screen_size.height as f32 / 16.0).ceil() as u32);
+        let workgroup_size = glam::UVec2::new((screen_size.width as f32 / 16.0).ceil() as u32,
+            (screen_size.height as f32 / 16.0).ceil() as u32);
         let light_grid = Self::create_grid(&device, workgroup_size);
         
         let depth_tex_pair = depth_texture.create_bind_group(&device);
@@ -205,7 +205,11 @@ impl LightCullCompute {
     pub fn update_screen_size(&mut self, size: PhysicalSize<u32>) {
         self.screen_size_buffer.write_buffer(&self.queue, 0,
             &[UVec2::new(size.width, size.height)]);
-        self.workgroup_size = glam::UVec2::new((size.width as f32 / 16.0).ceil() as u32, (size.height as f32 / 16.0).ceil() as u32);
+        self.workgroup_size = glam::UVec2::new((size.width as f32 / 16.0).ceil() as u32, 
+            (size.height as f32 / 16.0).ceil() as u32);
+
+        // I hate that the entire bind group is recreated on a resize but its the only way :(
+        self.light_indices_grid = Self::create_grid(&self.device, self.workgroup_size);
     }
 
     pub fn compute(&mut self, camera_buffers: &BufferWrapper, lights_buffers: &LightUniformBuffers, depth_texture: &RenderTexture) {
@@ -228,6 +232,7 @@ impl LightCullCompute {
             
             pass.dispatch_workgroups(self.workgroup_size.x, self.workgroup_size.y, 1);
         }
+
         self.queue.submit(std::iter::once(encoder.finish()));
         self.device.poll(wgpu::Maintain::Wait);
         self.cleanup();

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

@@ -238,5 +238,50 @@ fn blinn_phong_point_light(world_pos: vec3<f32>, world_norm: vec3<f32>, point_li
 }
 
 fn blinn_phong_spot_light(world_pos: vec3<f32>, world_norm: vec3<f32>, spot_light: Light, material: Material, specular_factor: vec3<f32>) -> vec3<f32> {
-    return vec3<f32>(0.0); // TODO
+    let light_color = spot_light.color;
+    let light_pos = spot_light.position;
+    let camera_view_pos = u_camera.position;
+
+    let light_dir = normalize(spot_light.position - world_pos);
+
+    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.spot_cutoff - spot_light.spot_outer_cutoff;
+    let intensity = clamp((theta - spot_light.spot_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 = calc_attenuation(spot_light, distance);
+
+    ambient_color *= attenuation * intensity;
+    diffuse_color *= attenuation * intensity;
+    specular_color *= attenuation * intensity;
+    //// end of spot light attenuation ////
+
+
+    return /*ambient_color +*/ diffuse_color + specular_color;
+}
+
+fn calc_attenuation(light: Light, distance: f32) -> f32 {
+    return 1.0 - smoothstep(light.range * light.smoothness, light.range, distance);
 }

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

@@ -33,6 +33,13 @@ struct Lights {
     data: array<Light>,
 };
 
+struct Cone {
+    tip: vec3<f32>,
+    height: f32,
+    direction: vec3<f32>,
+    radius: f32,
+}
+
 var<workgroup> wg_min_depth: atomic<u32>;
 var<workgroup> wg_max_depth: atomic<u32>;
 var<workgroup> wg_light_index_start: atomic<u32>;
@@ -160,13 +167,22 @@ fn cs_main(
                     && sphere_inside_frustrum(wg_frustum_planes, position, radius)) {
                 // TODO: add the light to the transparent geometry list
 
-                add_light(light_index);
-
                 if (!sphere_inside_plane(position, radius, wg_frustum_planes[4])) {
+                    add_light(light_index);
+                }
+            } else if (light.light_ty == LIGHT_TY_SPOT) {
+                let dir_vs = u_camera.view * vec4<f32>(light.direction, 1.0);
+                let cone_radius = tan(light.spot_cutoff) * light.range;
+                let cone = Cone(position, radius, dir_vs.xyz, cone_radius);
 
+                if (cone_inside_frustum(cone, wg_frustum_planes)) {
+                    // TODO: add the light to the transparent geometry list
+
+                    if (!cone_inside_plane(cone, wg_frustum_planes[4])) {
+                        add_light(light_index);
+                    }
                 }
             }
-            // TODO: spotlights
         }
     }
 
@@ -259,3 +275,44 @@ fn compute_plane(p0: vec3<f32>, p1: vec3<f32>, p2: vec3<f32>) -> vec4<f32> {
 
     return plane;
 }
+
+fn point_inside_plane(point: vec3<f32>, plane: vec4<f32>) -> bool {
+    return dot(plane.xyz, point) - plane.w < 0.0;
+}
+
+/// Check to see if a cone if fully behind (inside the negative halfspace of) a plane.
+///
+/// Source: Real-time collision detection, Christer Ericson (2005)
+///   (https://www.3dgep.com/forward-plus/#light-culling-compute-shader)
+fn cone_inside_plane(cone: Cone, plane: vec4<f32>) -> bool {
+    // Compute the farthest point on the end of the cone to the positive space of the plane.
+    let m = cross(cross(plane.xyz, cone.direction), cone.direction);
+    let farthest = cone.tip + cone.direction * cone.height - m * cone.radius;
+
+    // The cone is in the negative halfspace of the plane if the tip of the cone,
+    // and the farthest point on the end of the cone are inside the negative halfspace
+    // of the plane.
+    return point_inside_plane(cone.tip, plane) && point_inside_plane(farthest, plane);
+}
+
+fn cone_inside_frustum(cone: Cone, frustum: array<vec4<f32>, 6>) -> bool {
+    //let near_plane = frustum[4];
+    //let far_plane = frustum[5];
+
+    // check near and far clipping planes first
+    //if (cone_inside_plane(cone, near_plane) || cone_inside_plane(cone, far_plane)) {
+    //    return false;
+    //}
+
+    // to be able to index this array with a non-const value,
+    // it must be defined as a var
+    var frustum_v = frustum;
+
+    for (var i = 0u; i < 4u; i++) {
+        if (cone_inside_plane(cone, frustum_v[i])) {
+            return false;
+        }
+    }
+
+    return true;
+}

shell.nix

@@ -15,13 +15,19 @@ in
       mold
       udev
       lua5_4_compat
-      (nixpkgs.rustChannelOf { rustToolchain = ./rust-toolchain.toml; }).rust
+      ((nixpkgs.rustChannelOf { rustToolchain = ./rust-toolchain.toml; }).rust.override {
+        extensions = [
+          "rust-src"
+          "rust-analysis"
+        ];
+      })
     ];
     buildInputs = [
       udev alsa-lib libGL gcc
       vulkan-loader vulkan-headers vulkan-tools
       xorg.libX11 xorg.libXcursor xorg.libXi xorg.libXrandr # To use the x11 feature
       libxkbcommon wayland # To use the wayland feature
+      
     ];
     LD_LIBRARY_PATH = lib.makeLibraryPath buildInputs;
   }