move some stuff out of testbed into lyra-game

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

examples/testbed/Cargo.toml

@@ -13,4 +13,3 @@ lyra-engine = { path = "../../", version = "0.0.1" }
 anyhow = "1.0.75"
 async-std = "1.12.0"
 tracing = "0.1.37"
-fps_counter = "2.0.0"

examples/testbed/src/free_fly_camera.rs

@@ -1,102 +0,0 @@
-use lyra_engine::{
-    ecs::{query::{Res, View}, Component}, game::Game, input::ActionHandler, math::{EulerRot, Quat, Vec3}, plugin::Plugin, scene::CameraComponent, DeltaTime
-};
-
-/* enum FreeFlyCameraActions {
-    MoveForwardBackward,
-    MoveLeftRight,
-    MoveUpDown,
-    LookLeftRight,
-    LookUpDown,
-    LookRoll,
-} */
-
-#[derive(Clone, Component)]
-pub struct FreeFlyCamera {
-    pub speed: f32,
-    pub slow_speed_factor: f32,
-    pub look_speed: f32,
-    pub mouse_sensitivity: f32,
-    pub look_with_keys: bool,
-}
-
-impl Default for FreeFlyCamera {
-    fn default() -> Self {
-        Self {
-            speed: 4.0,
-            slow_speed_factor: 0.25,
-            look_speed: 0.3,
-            mouse_sensitivity: 1.0,
-            look_with_keys: false,
-        }
-    }
-}
-
-impl FreeFlyCamera {
-    #[allow(dead_code)]
-    pub fn new(speed: f32, slow_speed_factor: f32, look_speed: f32, mouse_sensitivity: f32, look_with_keys: bool) -> Self {
-        Self {
-            speed,
-            slow_speed_factor,
-            look_speed,
-            mouse_sensitivity,
-            look_with_keys,
-        }
-    }
-}
-
-pub fn free_fly_camera_controller(delta_time: Res<DeltaTime>, handler: Res<ActionHandler>, view: View<(&mut CameraComponent, &FreeFlyCamera)>) -> anyhow::Result<()> {
-    let delta_time = **delta_time;
-    for (mut cam, fly) in view.into_iter() {
-        let forward = cam.transform.forward();
-        let left = cam.transform.left();
-        let up = Vec3::Y;
-
-        let move_y = handler.get_axis_modifier("MoveUpDown").unwrap_or(0.0);
-        let move_x = handler.get_axis_modifier("MoveLeftRight").unwrap_or(0.0);
-        let move_z = handler.get_axis_modifier("MoveForwardBackward").unwrap_or(0.0);
-
-        let mut velocity = Vec3::ZERO;
-        velocity += move_y * up;
-        velocity += move_x * left;
-        velocity += move_z * forward;
-
-        if velocity != Vec3::ZERO {
-            cam.transform.translation += velocity.normalize() * fly.speed * delta_time; // TODO: speeding up
-        }
-
-        let motion_x = handler.get_axis_modifier("LookLeftRight").unwrap_or(0.0);
-        let motion_y = handler.get_axis_modifier("LookUpDown").unwrap_or(0.0);
-        let motion_z = handler.get_axis_modifier("LookRoll").unwrap_or(0.0);
-        
-        let mut camera_rot = Vec3::ZERO;
-        camera_rot.y -= motion_x * fly.mouse_sensitivity;
-        camera_rot.x -= motion_y * fly.mouse_sensitivity;
-        camera_rot.z -= motion_z * fly.mouse_sensitivity;
-
-        if camera_rot != Vec3::ZERO {
-            let look_velocity = camera_rot * fly.look_speed * delta_time;
-
-            let (mut y, mut x, _) = cam.transform.rotation.to_euler(EulerRot::YXZ);
-            x += look_velocity.x;
-            y += look_velocity.y;
-            x = x.clamp(-1.54, 1.54);
-
-            // rotation is not commutative, keep this order to avoid unintended roll
-            cam.transform.rotation = Quat::from_axis_angle(Vec3::Y, y)
-                * Quat::from_axis_angle(Vec3::X, x);
-        }
-    }
-
-    Ok(())
-}
-
-/// A plugin that adds the free fly camera controller system to the world. It is expected that
-/// there is a [`FreeFlyCamera`] in the world, if there isn't, the camera would not move.
-pub struct FreeFlyCameraPlugin;
-
-impl Plugin for FreeFlyCameraPlugin {
-    fn setup(&self, game: &mut Game) {
-        game.with_system("free_fly_camera_system", free_fly_camera_controller, &[]);
-    }
-}

examples/testbed/src/main.rs

@@ -1,28 +1,8 @@
 use std::ptr::NonNull;
 
-use lyra_engine::{assets::gltf::Gltf, ecs::{system::{BatchedSystem, Criteria, CriteriaSchedule, IntoSystem}, Component, World}, game::Game, input::{Action, ActionHandler, ActionKind, ActionMapping, ActionMappingId, ActionSource, InputActionPlugin, KeyCode, LayoutId, MouseAxis, MouseInput}, math::{self, Transform, Vec3}, render::light::{directional::DirectionalLight, SpotLight}, scene::CameraComponent, DeltaTime};
+use lyra_engine::{assets::gltf::Gltf, change_tracker::Ct, ecs::{query::{QueryBorrow, Res, View, ViewState}, system::{BatchedSystem, Criteria, CriteriaSchedule, IntoSystem}, Component, World}, game::Game, input::{Action, ActionHandler, ActionKind, ActionMapping, ActionMappingId, ActionSource, InputActionPlugin, KeyCode, LayoutId, MouseAxis, MouseInput}, math::{self, Quat, Transform, Vec3}, render::{light::{directional::DirectionalLight, PointLight, SpotLight}, window::{CursorGrabMode, WindowOptions}}, scene::{CameraComponent, FreeFlyCamera, FreeFlyCameraPlugin, ACTLBL_LOOK_LEFT_RIGHT, ACTLBL_LOOK_ROLL, ACTLBL_LOOK_UP_DOWN, ACTLBL_MOVE_FORWARD_BACKWARD, ACTLBL_MOVE_LEFT_RIGHT, ACTLBL_MOVE_UP_DOWN}, DeltaTime};
 use lyra_engine::assets::ResourceManager;
 
-mod free_fly_camera;
-use free_fly_camera::{FreeFlyCameraPlugin, FreeFlyCamera};
-
-#[derive(Clone, Copy, Hash, Debug)]
-pub enum ActionLabel {
-    MoveForwardBackward,
-    MoveLeftRight,
-    MoveUpDown,
-    LookLeftRight,
-    LookUpDown,
-    LookRoll,
-}
-
-const ACTLBL_MOVE_UP_DOWN: &str = "MoveUpDown";
-const ACTLBL_MOVE_LEFT_RIGHT: &str = "MoveLeftRight";
-const ACTLBL_MOVE_FORWARD_BACKWARD: &str = "MoveForwardBackward";
-const ACTLBL_LOOK_LEFT_RIGHT: &str = "LookLeftRight";
-const ACTLBL_LOOK_UP_DOWN: &str = "LookUpDown";
-const ACTLBL_LOOK_ROLL: &str = "LookRoll";
-
 struct FixedTimestep {
     max_tps: u32,
     fixed_time: f32,
@@ -93,15 +73,16 @@ async fn main() {
         //let diffuse_texture = resman.request::<Texture>("assets/happy-tree.png").unwrap();
         //let antique_camera_model = resman.request::<Model>("assets/AntiqueCamera.glb").unwrap();
         //let cube_model = resman.request::<Model>("assets/cube-texture-bin.glb").unwrap();
-        /* let cube_gltf = resman.request::<Gltf>("assets/texture-sep/texture-sep.gltf").unwrap();
+        let cube_gltf = resman.request::<Gltf>("assets/texture-sep/texture-sep.gltf").unwrap();
-        let crate_gltf = resman.request::<Gltf>("assets/crate/crate.gltf").unwrap();
+        /*let crate_gltf = resman.request::<Gltf>("assets/crate/crate.gltf").unwrap();
 
         let separate_gltf = resman.request::<Gltf>("assets/pos-testing/child-node-cubes.glb").unwrap(); */
         //drop(resman);
 
-        /* let cube_mesh = &cube_gltf.data_ref()
+        cube_gltf.wait_recurse_dependencies_load();
+        let cube_mesh = &cube_gltf.data_ref()
             .unwrap().meshes[0];
-        let crate_mesh = &crate_gltf.data_ref()
+        /* let crate_mesh = &crate_gltf.data_ref()
             .unwrap().meshes[0];
 
         let separate_scene = &separate_gltf.data_ref()
@@ -119,21 +100,6 @@ async fn main() {
             Transform::from_xyz(0.0, 0.0, 0.0),
         ));
 
-        /* world.spawn((
-            separate_scene.clone(),
-            Transform::from_xyz(0.0, -5.0, -10.0),
-        )); */
-
-        /* {
-            let cube_tran = Transform::from_xyz(-5.9026427, -1.8953488, -10.0);
-            //cube_tran.rotate_y(math::Angle::Degrees(180.0));
-            world.spawn((
-                cube_tran,
-                crate_mesh.clone(),
-                CubeFlag,
-            ));
-        } */
-
         {
             let mut light_tran = Transform::from_xyz(1.5, 2.5, 0.0);
             light_tran.scale = Vec3::new(0.5, 0.5, 0.5);
@@ -141,39 +107,76 @@ async fn main() {
             light_tran.rotate_y(math::Angle::Degrees(25.0));
             world.spawn((
                 DirectionalLight {
-                    color: Vec3::new(1.0, 1.0, 1.0),
+                    enabled: true,
-                    ambient: 0.3,
+                    color: Vec3::ONE,
-                    diffuse: 1.0,
+                    intensity: 0.35
-                    specular: 1.3,
+                    //..Default::default()
                 },
                 light_tran,
-                //cube_mesh.clone(),
-            ));
-        }
-
-        /* {
-            let mut light_tran = Transform::from_xyz(-3.5, 0.2, -4.5);
-            light_tran.scale = Vec3::new(0.5, 0.5, 0.5);
-            world.spawn((
-                SpotLight {
-                    color: Vec3::new(1.0, 0.2, 0.2),
-                    cutoff: math::Angle::Degrees(12.5),
-                    outer_cutoff: math::Angle::Degrees(17.5),
-
-                    constant: 1.0,
-                    linear: 0.007,
-                    quadratic: 0.0002,
-
-                    ambient: 0.0,
-                    diffuse: 7.0,
-                    specular: 1.0,
-                },
-                Transform::from(light_tran),
-                cube_mesh.clone(),
             ));
         }
 
         {
+            world.spawn((
+                PointLight {
+                    enabled: true,
+                    color: Vec3::new(0.0, 0.0, 1.0),
+                    intensity: 1.0,
+                    range: 2.0,
+                    ..Default::default()
+                },
+                Transform::new(
+                    //Vec3::new(-5.0, 1.0, -1.28), 
+                    Vec3::new(-5.0, 1.0, -0.0),
+                    //Vec3::new(-10.0, 0.94, -0.28), 
+
+                    Quat::IDENTITY,
+                    Vec3::new(0.25, 0.25, 0.25),
+                ),
+                cube_mesh.clone(),
+            ));
+
+            world.spawn((
+                PointLight {
+                    enabled: true,
+                    color: Vec3::new(0.0, 0.5, 1.0),
+                    intensity: 1.0,
+                    range: 1.0,
+                    ..Default::default()
+                },
+                Transform::new(
+                    Vec3::new(-3.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(),
+            ));
+
+            world.spawn((
+                SpotLight {
+                    enabled: true,
+                    color: Vec3::new(1.0, 0.0, 0.0),
+                    intensity: 1.0,
+                    range: 1.5,
+                    //cutoff: math::Angle::Degrees(45.0),
+                    ..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(),
+            ));
+        }
+
+        /* {
             let mut light_tran = Transform::from_xyz(2.0, 2.5, -9.5);
             light_tran.scale = Vec3::new(0.5, 0.5, 0.5);
             world.spawn((
@@ -218,54 +221,30 @@ async fn main() {
         } */
 
         let mut camera = CameraComponent::new_3d();
-        camera.transform.translation += math::Vec3::new(0.0, 0.0, 5.5);
+        // these values were taken by manually positioning the camera in the scene.
+        camera.transform = Transform::new(
+            Vec3::new(-10.0, 0.94, -0.28), 
+            Quat::from_xyzw(0.03375484, -0.7116095, 0.0342693, 0.70092666), 
+            Vec3::ONE
+        );
+        //camera.transform.translation += math::Vec3::new(0.0, 0.0, 5.5);
         world.spawn(( camera, FreeFlyCamera::default() ));
 
         Ok(())
     };
 
-    #[allow(unused_variables)]
+    let camera_debug_plugin = move |game: &mut Game| {
-    let fps_system = |world: &mut World| -> anyhow::Result<()> {
+        let sys = |handler: Res<ActionHandler>, view: View<&mut CameraComponent>| -> anyhow::Result<()> {
-        let mut counter = world.get_resource_mut::<fps_counter::FPSCounter>();
+            if handler.was_action_just_pressed("Debug") {
+                for mut cam in view.into_iter() {
+                    cam.debug = !cam.debug;
+                }
+            }
 
-        let fps = counter.tick();
+            Ok(())
+        };
 
-        println!("FPS: {fps}");
+        game.with_system("camera_debug_trigger", sys, &[]);
-
-        Ok(())
-    };
-    let _fps_plugin = move |game: &mut Game| {
-        let world = game.world_mut();
-        world.add_resource(fps_counter::FPSCounter::new());
-    };
-
-    let spin_system = |world: &mut World| -> anyhow::Result<()> {
-        const SPEED: f32 = 4.0;
-        let delta_time = **world.get_resource::<DeltaTime>();
-
-        for (mut transform, _) in world.view_iter::<(&mut Transform, &CubeFlag)>() {
-            let t = &mut transform;
-            t.rotate_y(math::Angle::Degrees(SPEED * delta_time));
-        }
-
-        for (mut transform, _s) in world.view_iter::<(&mut Transform, &mut SpotLight)>() {
-            let t = &mut transform;
-            t.rotate_x(math::Angle::Degrees(SPEED * delta_time));
-        }
-
-        Ok(())
-    };
-
-    let jiggle_plugin = move |game: &mut Game| {
-        game.world_mut().add_resource(TpsAccumulator(0.0));
-        
-        let mut sys = BatchedSystem::new();
-        sys.with_criteria(FixedTimestep::new(45));
-        sys.with_system(spin_system.into_system());
-        //sys.with_system(fps_system);
-
-        //game.with_system("fixed", sys, &[]);
-        //fps_plugin(game);
     };
 
     let action_handler_plugin = |game: &mut Game| {
@@ -278,6 +257,7 @@ async fn main() {
             .add_action(ACTLBL_LOOK_LEFT_RIGHT, Action::new(ActionKind::Axis))
             .add_action(ACTLBL_LOOK_UP_DOWN, Action::new(ActionKind::Axis))
             .add_action(ACTLBL_LOOK_ROLL, Action::new(ActionKind::Axis))
+            .add_action("Debug", Action::new(ActionKind::Button))
 
             .add_mapping(ActionMapping::builder(LayoutId::from(0), ActionMappingId::from(0))
                 .bind(ACTLBL_MOVE_FORWARD_BACKWARD, &[
@@ -308,6 +288,9 @@ async fn main() {
                     ActionSource::Keyboard(KeyCode::E).into_binding_modifier(-1.0),
                     ActionSource::Keyboard(KeyCode::Q).into_binding_modifier(1.0),
                 ])
+                .bind("Debug", &[
+                    ActionSource::Keyboard(KeyCode::B).into_binding(),
+                ])
                 .finish()
             ).finish();
 
@@ -337,7 +320,7 @@ async fn main() {
         .with_plugin(action_handler_plugin)
         //.with_plugin(script_test_plugin)
         //.with_plugin(fps_plugin)
-        .with_plugin(jiggle_plugin)
+        .with_plugin(camera_debug_plugin)
         .with_plugin(FreeFlyCameraPlugin)
         .run().await;
 }

examples/testbed/testbed-renderdoc.cap

@@ -5,7 +5,7 @@
         "commandLine": "",
         "environment": [
         ],
-        "executable": "/media/data_drive/Development/Rust/lyra-test/engine/target/debug/testbed",
+        "executable": "/media/data_drive/Development/Rust/lyra-engine/target/debug/testbed",
         "inject": false,
         "numQueuedFrames": 1,
         "options": {
@@ -23,6 +23,6 @@
             "verifyBufferAccess": false
         },
         "queuedFrameCap": 5,
-        "workingDir": "/media/data_drive/Development/Rust/lyra-test/engine/examples/testbed"
+        "workingDir": "/media/data_drive/Development/Rust/lyra-engine/examples/testbed"
     }
 }

lyra-game/src/game.rs

@@ -350,6 +350,7 @@ impl Game {
                 // done by prefix, so it includes all lyra subpackages
                 .with_target("lyra", Level::DEBUG)
                 .with_target("wgpu", Level::WARN)
+                .with_target("winit", Level::DEBUG)
                 .with_default(Level::INFO))
             .init();
 

lyra-game/src/input/buttons.rs

@@ -93,11 +93,24 @@ impl<T: Button> InputButtons<T> {
         let hash = Self::get_button_hash(&button);
         match self.button_events.get(&hash) {
             Some(button_event) => match button_event {
-                // this if statement should always be true, but just in case ;)
                 ButtonEvent::JustPressed(b) if button == *b => true,
                 _ => false,
             },
             None => false
         }
     }
+
+    /// Update any JustPressed events into Pressed events
+    /// 
+    /// This must be done so that a key does not stay as JustPressed between multiple ticks
+    pub fn update(&mut self) {
+        for bev in self.button_events.values_mut() {
+            match bev {
+                ButtonEvent::JustPressed(btn) => {
+                    *bev = ButtonEvent::Pressed(btn.clone());
+                },
+                _ => {},
+            }
+        }
+    }
 }

lyra-game/src/input/system.rs

@@ -102,12 +102,15 @@ impl crate::ecs::system::System for InputSystem {
         let queue = world.try_get_resource_mut::<EventQueue>()
             .and_then(|q| q.read_events::<InputEvent>());
 
+        let mut e = world.get_resource_or_else(InputButtons::<winit::event::VirtualKeyCode>::new);
+        e.update();
+        drop(e);
+
         if queue.is_none() {
             return Ok(());
         }
 
         let mut events = queue.unwrap();
-        
         while let Some(event) = events.pop_front() {
             self.process_event(world, &event);
         }

lyra-game/src/render/camera.rs

@@ -38,16 +38,43 @@ impl Projection {
 #[repr(C)]
 #[derive(Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)]
 pub struct CameraUniform {
-    pub view_proj: glam::Mat4,
+    /// The view matrix of the camera
-    // vec4 is used because of the uniforms 16 byte spacing requirement
+    pub view: glam::Mat4,
-    pub view_pos: glam::Vec4, 
+    /// The inverse of the projection matrix of the camera
+    pub inverse_projection: glam::Mat4,
+    /// The view projection matrix 
+    pub view_projection: glam::Mat4,
+    pub projection: glam::Mat4,
+    /// The position of the camera
+    pub position: glam::Vec3,
+    pub tile_debug: u32,
+    //_padding: [u8; 3],
+
 }
 
 impl Default for CameraUniform {
     fn default() -> Self {
         Self {
-            view_proj: glam::Mat4::IDENTITY,
+            view: glam::Mat4::IDENTITY,
-            view_pos: Default::default()
+            inverse_projection: glam::Mat4::IDENTITY,
+            view_projection: glam::Mat4::IDENTITY,
+            projection: glam::Mat4::IDENTITY,
+            position: Default::default(),
+            tile_debug: 0,
+            //_padding: 0,
+        }
+    }
+}
+
+impl CameraUniform {
+    pub fn new(view: glam::Mat4, inverse_projection: glam::Mat4, view_projection: glam::Mat4, projection: glam::Mat4, position: glam::Vec3) -> Self {
+        Self {
+            view,
+            inverse_projection,
+            view_projection,
+            projection,
+            position,
+            tile_debug: 0
         }
     }
 }
@@ -79,26 +106,38 @@ impl RenderCamera {
         self.aspect = size.width as f32 / size.height as f32;
     }
 
-    pub fn update_view_projection(&mut self, camera: &CameraComponent) -> &glam::Mat4 {
+    /// Calculates the view projection, and the view
+    /// 
+    /// Returns: A tuple with the view projection as the first element, and the
+    /// view matrix as the second.
+    pub fn calc_view_projection(&mut self, camera: &CameraComponent) -> CameraUniform {
+        let position = camera.transform.translation;
+        let forward = camera.transform.forward();
+        let up = camera.transform.up();
+
+        let view = glam::Mat4::look_to_rh(
+            position,
+            forward,
+            up
+        );
+
         match camera.mode {
             CameraProjectionMode::Perspective => {
-                let position = camera.transform.translation;
-                let forward = camera.transform.forward();
-                let up = camera.transform.up();
-
-                let view = glam::Mat4::look_to_rh(
-                    position,
-                    forward,
-                    up
-                );
-                
                 let proj = glam::Mat4::perspective_rh_gl(camera.fov.to_radians(), self.aspect, self.znear, self.zfar);
 
                 self.view_proj = OPENGL_TO_WGPU_MATRIX * proj * view;
-                &self.view_proj
+                //(&self.view_proj, view)
+
+                CameraUniform {
+                    view,
+                    inverse_projection: proj.inverse(),
+                    view_projection: self.view_proj,
+                    projection: proj,
+                    position,
+                    tile_debug: camera.debug as u32,
+                }
             },
             CameraProjectionMode::Orthographic => {
-                let position = camera.transform.translation;
                 let target = camera.transform.rotation * glam::Vec3::new(0.0, 0.0, -1.0);
                 let target = target.normalize();
 
@@ -111,7 +150,15 @@ impl RenderCamera {
                 let proj = glam::Mat4::orthographic_rh_gl(-size_x, size_x, -size_y, size_y, self.znear, self.zfar);
 
                 self.view_proj = OPENGL_TO_WGPU_MATRIX * proj;
-                &self.view_proj
+                
+                CameraUniform {
+                    view,
+                    inverse_projection: proj.inverse(),
+                    view_projection: self.view_proj,
+                    projection: proj,
+                    position,
+                    tile_debug: camera.debug as u32,
+                }
             },
         }
     }

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

@@ -1,11 +1,18 @@
 use lyra_ecs::Component;
 
-#[derive(Default, Debug, Clone, Component)]
+#[derive(Debug, Clone, Component)]
 pub struct DirectionalLight {
-    //pub direction: glam::Quat,
+    pub enabled: bool,
     pub color: glam::Vec3,
-
+    pub intensity: f32,
-    pub ambient: f32,
+}
-    pub diffuse: f32,
+
-    pub specular: f32,
+impl Default for DirectionalLight {
+    fn default() -> Self {
+        Self {
+            enabled: true,
+            color: glam::Vec3::new(1.0, 1.0, 1.0),
+            intensity: 1.0,
+        }
+    }
 }

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

@@ -6,14 +6,14 @@ use lyra_ecs::{Entity, Tick, World, query::{Entities, TickOf}};
 pub use point::*;
 pub use spotlight::*;
 
-use std::{collections::{VecDeque, HashMap}, marker::PhantomData};
+use std::{collections::{HashMap, VecDeque}, marker::PhantomData, mem};
-
-use std::mem;
 
 use crate::math::Transform;
 
 use self::directional::DirectionalLight;
 
+use super::render_buffer::BindGroupPair;
+
 const MAX_LIGHT_COUNT: usize = 16;
 
 /// A struct that stores a list of lights in a wgpu::Buffer.
@@ -101,20 +101,23 @@ impl<U: Default + bytemuck::Pod + bytemuck::Zeroable> LightBuffer<U> {
 
 pub(crate) struct LightUniformBuffers {
     pub buffer: wgpu::Buffer,
-    pub bindgroup_layout: wgpu::BindGroupLayout,
+    pub bind_group_pair: BindGroupPair,
-    pub bindgroup: wgpu::BindGroup,
+    pub light_indexes: HashMap<Entity, u32>,
-    pub lights_uniform: LightsUniform,
+    dead_indices: VecDeque<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: 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(
             &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,209 +152,181 @@ 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,
+            bind_group_pair: BindGroupPair::new(bindgroup, bindgroup_layout),
-            bindgroup,
+            light_indexes: Default::default(),
-            lights_uniform: LightsUniform::default(),
+            current_light_idx: 0,
-            point_lights,
+            dead_indices: VecDeque::new(),
-            spot_lights,
         }
     }
 
+    /// 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) {
+        // 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>)>() {
 
-            if !self.point_lights.has_light(entity) || light_epoch == world_tick || transform_epoch == world_tick {
+            let (new, idx) = self.get_index_for(&mut missed_lights, entity);
-                let uniform = PointLightUniform::from_bundle(&point_light, &transform);
+            self.light_indexes.insert(entity, idx);
-                self.point_lights.update_or_add(&mut self.lights_uniform.point_lights, entity, uniform);
+
-                //debug!("Updated point light");
+            if new || light_epoch == world_tick || transform_epoch == world_tick {
+                let uniform = LightUniform::from_point_light_bundle(&point_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, spot_light, transform, light_epoch, transform_epoch)
                 in world.view_iter::<(Entities, &SpotLight, &Transform, TickOf<SpotLight>, TickOf<Transform>)>() {
 
-            if !self.spot_lights.has_light(entity) || light_epoch == world_tick || transform_epoch == world_tick {
+            let (new, idx) = self.get_index_for(&mut missed_lights, entity);
-                let uniform = SpotLightUniform::from_bundle(&spot_light, &transform);
+            self.light_indexes.insert(entity, idx);
-                self.spot_lights.update_or_add(&mut self.lights_uniform.spot_lights, entity, uniform);
+
-                //debug!("Updated spot light");
+            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]));
             }
         }
 
-        if let Some((dir_light, transform)) = 
+        for (entity, dir_light, transform, light_epoch, transform_epoch)
-                world.view_iter::<(&DirectionalLight, &Transform)>().next() {
+                in world.view_iter::<(Entities, &DirectionalLight, &Transform, TickOf<DirectionalLight>, TickOf<Transform>)>() {
             
-            let uniform = DirectionalLightUniform::from_bundle(&dir_light, &transform);
+            let (new, idx) = self.get_index_for(&mut missed_lights, entity);
-            self.lights_uniform.directional_light = uniform;
+            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);
+                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]));
+            }
         }
 
-        self.lights_uniform.point_light_count = self.point_lights.buffer_count as u32;
+        // anything left in missed_lights were lights that were deleted
-        self.lights_uniform.spot_light_count = self.spot_lights.buffer_count as u32;
+        let len = missed_lights.len();
-        queue.write_buffer(&self.buffer, 0, bytemuck::cast_slice(&[self.lights_uniform]));
+        self.dead_indices.reserve(len);
+
+        for (_, v) in missed_lights.drain() {
+            // 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]));
     }
 }
 
 #[repr(C)]
-#[derive(Default, Debug, Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)]
+#[derive(Default, Debug, Copy, Clone)]
-pub struct LightsUniform {
+pub(crate) enum LightType {
-    point_lights: [PointLightUniform; MAX_LIGHT_COUNT],
+    #[default]
-    point_light_count: u32,
+    Directional = 0,
-    _padding: [u32; 3],
+    Point = 1,
-    spot_lights: [SpotLightUniform; MAX_LIGHT_COUNT],
+    Spotlight = 2,
-    spot_light_count: u32,
-    _padding2: [u32; 3],
-    directional_light: DirectionalLightUniform,
 }
 
 #[repr(C)]
 #[derive(Default, Debug, Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)]
-pub struct PointLightUniform {
+pub(crate) struct LightUniform {
-    /// 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,
-            linear: light.linear,
-            quadratic: light.quadratic,
-
-            ambient: light.ambient,
-            diffuse: light.diffuse,
-            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 light_type: u32, // enum LightType
     pub direction: glam::Vec3,
-    pub _padding2: u32,
+    pub enabled: u32, // bool
     pub color: glam::Vec3,
-    // no padding is needed here since cutoff acts as the padding
+    // no padding is needed here since range acts as the padding
     // that would usually be needed for the vec3
 
-    pub cutoff: f32,
+    pub range: f32,
-    pub outer_cutoff: f32,
+    pub intensity: f32,
+    pub smoothness: f32,
 
-    /// The constant used in the quadratic attenuation calculation. Its best to leave this at 1.0
+    pub spot_cutoff_rad: f32,
-    pub constant: f32,
+    pub spot_outer_cutoff_rad: 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 {
+impl LightUniform {
-    /// Create the SpotLightUniform from an ECS bundle
+    pub fn from_point_light_bundle(light: &PointLight, transform: &Transform) -> Self {
-    pub fn from_bundle(light: &SpotLight, transform: &Transform) -> Self {
         Self {
+            light_type: LightType::Point 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(),
+            range: light.range,
-            constant: light.constant,
+            intensity: light.intensity,
-            linear: light.linear,
+            smoothness: light.smoothness,
-            quadratic: light.quadratic,
+
-            ambient: light.ambient,
+            spot_cutoff_rad: 0.0,
-            diffuse: light.diffuse,
+            spot_outer_cutoff_rad: 0.0,
-            specular: light.specular,
+
-            _padding3: 0,
+        }
+    }
+
+    pub fn from_directional_bundle(light: &DirectionalLight, transform: &Transform) -> Self {
+        Self {
+            light_type: LightType::Directional as u32,
+            enabled: light.enabled as u32,
+            position: transform.translation,
+            direction: transform.forward(),
+            color: light.color,
+
+            range: 0.0,
+            intensity: light.intensity,
+            smoothness: 0.0,
+
+            spot_cutoff_rad: 0.0,
+            spot_outer_cutoff_rad: 0.0,
+        }
+    }
+
+    // Create the SpotLightUniform from an ECS bundle
+    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,
+            direction: transform.forward(),
+            color: light.color,
+
+            range: light.range,
+            intensity: light.intensity,
+            smoothness: light.smoothness,
+
+            spot_cutoff_rad: light.cutoff.to_radians(),
+            spot_outer_cutoff_rad: light.outer_cutoff.to_radians(),
         }
     }
 }
+

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

@@ -1,13 +1,22 @@
 use lyra_ecs::Component;
 
-#[derive(Default, Debug, Clone, Component)]
+#[derive(Debug, Clone, Component)]
 pub struct PointLight {
+    pub enabled: bool,
     pub color: glam::Vec3,
+    pub range: f32,
     pub intensity: f32,
-    pub constant: f32,
+    pub smoothness: f32,
-    pub linear: f32,
+}
-    pub quadratic: f32,
+
-    pub ambient: f32,
+impl Default for PointLight {
-    pub diffuse: f32,
+    fn default() -> Self {
-    pub specular: f32,
+        Self {
+            enabled: true,
+            color: glam::Vec3::new(1.0, 1.0, 1.0),
+            range: 1.0,
+            intensity: 1.0,
+            smoothness: 0.75,
+        }
+    }
 }

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

@@ -0,0 +1,244 @@
+use std::{borrow::Cow, mem, rc::Rc};
+
+use glam::UVec2;
+use wgpu::{util::DeviceExt, ComputePipeline};
+use winit::dpi::PhysicalSize;
+
+use super::{light::LightUniformBuffers, render_buffer::{BindGroupPair, BufferWrapper}, texture::RenderTexture};
+
+#[allow(dead_code)]
+pub(crate) struct LightIndicesGridBuffer {
+    index_counter_buffer: wgpu::Buffer,
+    indices_buffer: wgpu::Buffer,
+    grid_texture: wgpu::Texture,
+    grid_texture_view: wgpu::TextureView,
+    pub bg_pair: BindGroupPair,
+}
+
+pub(crate) struct LightCullCompute {
+    device: Rc<wgpu::Device>,
+    queue: Rc<wgpu::Queue>,
+    pipeline: ComputePipeline,
+    pub light_indices_grid: LightIndicesGridBuffer,
+    screen_size_buffer: BufferWrapper,
+    workgroup_size: glam::UVec2,
+}
+
+impl LightCullCompute {
+    /// Create the LightIndiciesGridBuffer object
+    fn create_grid(device: &wgpu::Device, workgroup_size: glam::UVec2) -> LightIndicesGridBuffer {
+        let mut contents = Vec::<u8>::new();
+        let contents_len = workgroup_size.x * workgroup_size.y * 200 * mem::size_of::<u32>() as u32;
+        contents.resize(contents_len as _, 0);
+
+        let light_indices_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
+            label: Some("B_LightIndices"),
+            contents: &contents,
+            usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
+        });
+
+        let light_index_counter_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
+            label: Some("B_LightIndexCounter"),
+            contents: &bytemuck::cast_slice(&[0]),
+            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,
+                },
+                wgpu::BindGroupLayoutEntry {
+                    binding: 2,
+                    visibility: wgpu::ShaderStages::COMPUTE,
+                    ty: wgpu::BindingType::Buffer {
+                        ty: wgpu::BufferBindingType::Storage {
+                            read_only: false
+                        },
+                        has_dynamic_offset: false,
+                        min_binding_size: None,
+                    },
+                    count: None,
+                }
+            ],
+            label: Some("BGL_LightIndicesGrid"),
+        });
+
+        let size = wgpu::Extent3d {
+            width: workgroup_size.x,
+            height: workgroup_size.y,
+            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)
+                },
+                wgpu::BindGroupEntry {
+                    binding: 2,
+                    resource: wgpu::BindingResource::Buffer(
+                            wgpu::BufferBinding {
+                            buffer: &light_index_counter_buffer,
+                            offset: 0,
+                            size: None, // the entire light buffer is needed
+                        }
+                    )
+                },
+            ],
+            label: Some("BG_LightIndicesGrid"),
+        });
+
+        LightIndicesGridBuffer {
+            index_counter_buffer: light_index_counter_buffer,
+            indices_buffer: light_indices_buffer,
+            grid_texture,
+            grid_texture_view,
+            bg_pair: BindGroupPair::new(light_indices_bg, light_indices_bg_layout),
+        }
+    }
+
+    pub fn new(device: Rc<wgpu::Device>, queue: Rc<wgpu::Queue>, screen_size: PhysicalSize<u32>, lights_buffers: &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 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);
+
+        let layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+            label: Some("PipeLay_LightCull"),
+            bind_group_layouts: &[
+                &depth_tex_pair.layout,
+                &camera_buffers.bindgroup_layout().unwrap(),
+                &lights_buffers.bind_group_pair.layout,
+                &light_grid.bg_pair.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,
+            light_indices_grid: light_grid,
+            screen_size_buffer,
+            workgroup_size,
+        }
+    }
+
+    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);
+
+        // 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) {
+        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);
+
+            pass.set_bind_group(0, depth_texture.bind_group(), &[]);
+            pass.set_bind_group(1, &camera_buffers.bindgroup(), &[]);
+            pass.set_bind_group(2, &lights_buffers.bind_group_pair.bindgroup, &[]);
+            pass.set_bind_group(3, &self.light_indices_grid.bg_pair.bindgroup, &[]);
+            pass.set_bind_group(4, self.screen_size_buffer.bindgroup(), &[]);
+            
+            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();
+    }
+
+    pub fn cleanup(&mut self) {
+        self.queue.write_buffer(&self.light_indices_grid.index_counter_buffer, 0, &bytemuck::cast_slice(&[0]));
+    }
+}

lyra-game/src/render/material.rs

@@ -1,4 +1,4 @@
-use std::sync::Arc;
+use std::rc::Rc;
 
 use lyra_resource::{ResHandle, Texture};
 
@@ -11,7 +11,7 @@ pub struct MaterialSpecular {
     pub color_texture: Option<RenderTexture>,
 }
 
-fn texture_to_render(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: &Arc<wgpu::BindGroupLayout>, i: &Option<ResHandle<Texture>>) -> Option<RenderTexture> {
+fn texture_to_render(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: &Rc<wgpu::BindGroupLayout>, i: &Option<ResHandle<Texture>>) -> Option<RenderTexture> {
     if let Some(tex) = i {
         RenderTexture::from_resource(device, queue, bg_layout.clone(), tex, None).ok()
     } else {
@@ -20,7 +20,7 @@ fn texture_to_render(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: &Arc
 }
 
 impl MaterialSpecular {
-    pub fn from_resource(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Arc<wgpu::BindGroupLayout>, value: &lyra_resource::gltf::Specular) -> Self {
+    pub fn from_resource(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Rc<wgpu::BindGroupLayout>, value: &lyra_resource::gltf::Specular) -> Self {
         let tex  = texture_to_render(device, queue, &bg_layout, &value.texture);
         let color_tex  = texture_to_render(device, queue, &bg_layout, &value.color_texture);
 
@@ -45,7 +45,7 @@ pub struct Material {
 }
 
 impl Material {
-    pub fn from_resource(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Arc<wgpu::BindGroupLayout>, value: &lyra_resource::gltf::Material) -> Self {
+    pub fn from_resource(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Rc<wgpu::BindGroupLayout>, value: &lyra_resource::gltf::Material) -> Self {
         let diffuse_texture  = texture_to_render(device, queue, &bg_layout, &value.base_color_texture);
 
         let specular = value.specular.as_ref().map(|s| MaterialSpecular::from_resource(device, queue, bg_layout.clone(), s));

lyra-game/src/render/mod.rs

@@ -12,3 +12,4 @@ pub mod camera;
 pub mod window;
 pub mod transform_buffer_storage;
 pub mod light;
+pub mod light_cull_compute;

lyra-game/src/render/render_buffer.rs

@@ -1,4 +1,4 @@
-use std::{sync::Arc, num::NonZeroU32};
+use std::{num::NonZeroU32, rc::Rc};
 
 use wgpu::util::DeviceExt;
 
@@ -23,13 +23,13 @@ impl RenderBuffer {
 
 pub struct BindGroupPair {
     pub bindgroup: wgpu::BindGroup,
-    pub layout: Arc<wgpu::BindGroupLayout>,
+    pub layout: Rc<wgpu::BindGroupLayout>,
 }
 
 impl BindGroupPair {
-    pub fn new_from_layout(device: &wgpu::Device, layout: Arc<wgpu::BindGroupLayout>, entries: &[wgpu::BindGroupEntry<'_>]) -> Self {
+    pub fn create_bind_group(device: &wgpu::Device, layout: Rc<wgpu::BindGroupLayout>, entries: &[wgpu::BindGroupEntry<'_>]) -> Self {
         let bindgroup = device.create_bind_group(&wgpu::BindGroupDescriptor {
-            layout: layout.as_ref(),
+            layout: &layout,
             entries,
             label: None,
         });
@@ -39,12 +39,19 @@ impl BindGroupPair {
             layout,
         }
     }
+
+    pub fn new(bindgroup: wgpu::BindGroup, layout: wgpu::BindGroupLayout) -> Self {
+        Self {
+            bindgroup,
+            layout: Rc::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,15 +71,82 @@ 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,
+        }
+    }
+
+    /// Creates a builder for a BufferWrapper
     pub fn builder() -> BufferWrapperBuilder {
         BufferWrapperBuilder::new()
     }
+
+    /// Retrieve the layout of the bindgroup associated with this buffer.
+    /// 
+    /// Returns None if this buffer object was not provided a bindgroup.
+    pub fn bindgroup_layout(&self) -> Option<&wgpu::BindGroupLayout> {
+        self.bindgroup_pair.as_ref().map(|bg| &*bg.layout)
+    }
+
+    /// Queue's the data to be written to `buffer` starting at `offset`.
+    /// 
+    /// The write is not immediately submitted, and instead enqueued
+    /// internally to happen at the start of the next submit() call.
+    /// 
+    /// This method fails if data overruns the size of buffer starting at offset.
+    /// 
+    /// See [`wgpu::Queue::write_buffer`](https://docs.rs/wgpu/latest/wgpu/struct.Queue.html#method.write_buffer).
+    pub fn write_buffer<T>(&self, queue: &wgpu::Queue, offset: u64, data: &[T])
+    where
+        T: bytemuck::NoUninit
+    {
+        queue.write_buffer(&self.inner_buf, offset, bytemuck::cast_slice(data));
+    }
+
+    /// Sets the buffer's bind group to `index` in the `pass`.
+    /// 
+    /// The bind group layout in the active pipeline when any `draw()` function is called must
+    /// 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 render_pass_bind_at<'a, 'b>(
+        &'a self,
+        pass: &'b mut wgpu::RenderPass<'a>,
+        index: u32,
+        offsets: &[wgpu::DynamicOffset],
+    ) {
+        let pair = self.bindgroup_pair.as_ref().expect(
+            "BufferWrapper is missing bindgroup pair! Cannot set bind group on RenderPass!",
+        );
+        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
+/// 
+/// ```nobuild
+/// let camera_buffer = BufferWrapper::builder()
+///     .buffer_usage(wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST)
+///     .contents(&[CameraUniform::default()])
+///     .label_prefix("Camera")
+///     .visibility(wgpu::ShaderStages::VERTEX | wgpu::ShaderStages::FRAGMENT)
+///     .buffer_dynamic_offset(false)
+///     .finish(&device);
+/// ```
 #[derive(Default)]
 pub struct BufferWrapperBuilder {
     buffer_usage: Option<wgpu::BufferUsages>,
@@ -198,7 +272,7 @@ impl BufferWrapperBuilder {
                     ],
                     label: self.format_label("BGL_").as_deref(),
                 });
-                let bg_layout = Arc::new(bg_layout);
+                //let bg_layout = Arc::new(bg_layout);
                 
     
                 let bg = device.create_bind_group(&wgpu::BindGroupDescriptor {
@@ -214,7 +288,7 @@ impl BufferWrapperBuilder {
 
                 BindGroupPair {
                     bindgroup: bg,
-                    layout: bg_layout,
+                    layout: Rc::new(bg_layout),
                 }
             }
         };
@@ -222,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,13 +47,19 @@ 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)>,
 
     //#[allow(dead_code)]
     //render_texture: Option<RenderTexture>,
-    material: Option<Material>,
+    material: Option<Rc<Material>>,
 
     // The index of the transform for this entity.
     // The tuple is structured like this: (transform index, index of transform inside the buffer)
@@ -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>,
@@ -80,6 +88,7 @@ pub struct BasicRenderer {
     pub render_jobs: VecDeque<RenderJob>,
 
     mesh_buffers: HashMap<uuid::Uuid, MeshBufferStorage>, // TODO: clean up left over buffers from deleted entities/components
+    material_buffers: HashMap<uuid::Uuid, Rc<Material>>,
     entity_meshes: HashMap<Entity, uuid::Uuid>,
     entity_last_transforms: HashMap<Entity, CachedTransform>,
 
@@ -88,16 +97,18 @@ pub struct BasicRenderer {
     render_limits: Limits,
 
     inuse_camera: RenderCamera,
-    camera_buffer: wgpu::Buffer,
+    camera_buffer: BufferWrapper,
-    camera_bind_group: wgpu::BindGroup,
+    //camera_bind_group: wgpu::BindGroup,
 
-    bgl_texture: Arc<BindGroupLayout>,
+    bgl_texture: Rc<BindGroupLayout>,
     default_texture: RenderTexture,
     depth_buffer_texture: RenderTexture,
 
     material_buffer: BufferWrapper,
 
     light_buffers: LightUniformBuffers,
+
+    light_cull_compute: LightCullCompute,
 }
 
 impl BasicRenderer {
@@ -123,7 +134,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
@@ -143,10 +154,7 @@ impl BasicRenderer {
         let render_limits = device.limits();
         let surface_caps = surface.get_capabilities(&adapter);
 
-        let present_mode = surface_caps.present_modes[0]; /* match surface_caps.present_modes.contains(&wgpu::PresentMode::Immediate) {
+        let present_mode = surface_caps.present_modes[0];
-            true => wgpu::PresentMode::Immediate,
-            false => surface_caps.present_modes[0]
-        }; */
 
         debug!("present mode: {:?}", present_mode);
 
@@ -165,7 +173,7 @@ impl BasicRenderer {
         };
         surface.configure(&device, &config);
 
-        let bgl_texture = Arc::new(RenderTexture::create_layout(&device));
+        let bgl_texture = Rc::new(RenderTexture::create_layout(&device));
 
         let shader_src = include_str!("shaders/base.wgsl");
         let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -174,43 +182,15 @@ impl BasicRenderer {
         });
 
         let transform_buffers = TransformBuffers::new(&device);
+        let camera_buffer = BufferWrapper::builder()
+            .buffer_usage(wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST)
+            .contents(&[CameraUniform::default()])
+            .label_prefix("Camera")
+            .visibility(wgpu::ShaderStages::VERTEX | wgpu::ShaderStages::FRAGMENT | wgpu::ShaderStages::COMPUTE)
+            .buffer_dynamic_offset(false)
+            .finish(&device);
 
-        let camera_buffer = device.create_buffer_init(
+        let mut depth_texture = RenderTexture::create_depth_texture(&device, &config, "Tex_Depth");
-            &wgpu::util::BufferInitDescriptor {
-                label: Some("Camera Buffer"),
-                contents: bytemuck::cast_slice(&[CameraUniform::default()]),
-                usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
-            }
-        );
-
-        let camera_bind_group_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
-            entries: &[
-                wgpu::BindGroupLayoutEntry {
-                    binding: 0,
-                    visibility: wgpu::ShaderStages::VERTEX | wgpu::ShaderStages::FRAGMENT,
-                    ty: wgpu::BindingType::Buffer {
-                        ty: wgpu::BufferBindingType::Uniform,
-                        has_dynamic_offset: false,
-                        min_binding_size: None,
-                    },
-                    count: None,
-                }
-            ],
-            label: Some("camera_bind_group_layout"),
-        });
-        
-        let camera_bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
-            layout: &camera_bind_group_layout,
-            entries: &[
-                wgpu::BindGroupEntry {
-                    binding: 0,
-                    resource: camera_buffer.as_entire_binding(),
-                }
-            ],
-            label: Some("camera_bind_group"),
-        });
-
-        let depth_texture = RenderTexture::create_depth_texture(&device, &config, "Depth Buffer");
 
         // load the default texture
         let bytes = include_bytes!("default_texture.png");
@@ -222,9 +202,12 @@ impl BasicRenderer {
             .buffer_usage(wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST)
             .visibility(wgpu::ShaderStages::FRAGMENT)
             .contents(&[MaterialUniform::default()])
-            //.size(mem::size_of::<MaterialUniform>())
             .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,
@@ -241,6 +224,7 @@ impl BasicRenderer {
             render_pipelines: HashMap::new(),
             render_jobs: VecDeque::new(),
             mesh_buffers: HashMap::new(),
+            material_buffers: HashMap::new(),
             entity_meshes: HashMap::new(),
 
             render_limits,
@@ -248,7 +232,6 @@ impl BasicRenderer {
 
             inuse_camera: RenderCamera::new(size),
             camera_buffer,
-            camera_bind_group,
 
             bgl_texture,
             default_texture,
@@ -257,15 +240,19 @@ impl BasicRenderer {
 
             light_buffers: light_uniform_buffers,
             material_buffer: mat_buffer,
+            light_cull_compute,
         };
 
         // create the default pipelines
         let mut pipelines = HashMap::new();
         pipelines.insert(0, Arc::new(FullRenderPipeline::new(&s.device, &s.config, &shader, 
             vec![super::vertex::Vertex::desc(),],
-            vec![&s.bgl_texture, &s.transform_buffers.bindgroup_layout, &camera_bind_group_layout,
+            vec![&s.bgl_texture, &s.transform_buffers.bindgroup_layout, 
-            &s.light_buffers.bindgroup_layout, &s.material_buffer.bindgroup_pair.as_ref().unwrap().layout,
+            s.camera_buffer.bindgroup_layout().unwrap(),
-            &s.bgl_texture])));
+            &s.light_buffers.bind_group_pair.layout, &s.material_buffer.bindgroup_pair.as_ref().unwrap().layout,
+            &s.bgl_texture,
+            &s.light_cull_compute.light_indices_grid.bg_pair.layout,
+        ])));
         s.render_pipelines = pipelines;
 
         s
@@ -364,18 +351,24 @@ impl BasicRenderer {
         let (vertex_buffer, buffer_indices) = self.create_vertex_index_buffers(mesh);
 
         let material = mesh.material.as_ref()
-            .expect("Material resource not loaded yet")
+            .expect("Material resource not loaded yet");
-            .data_ref()
+        let material_ref = material.data_ref()
             .unwrap();
-        let material = Material::from_resource(&self.device, &self.queue, self.bgl_texture.clone(), &material);
+        
-        let uni = MaterialUniform::from(&material);
+        let material = self.material_buffers.entry(material.uuid())
+            .or_insert_with(|| {
+                debug!(uuid=material.uuid().to_string(), "Sending material to gpu");
+                Rc::new(Material::from_resource(&self.device, &self.queue, self.bgl_texture.clone(), &material_ref))
+            });
+
+        // TODO: support material uniforms from multiple uniforms
+        let uni = MaterialUniform::from(&**material);
         self.queue.write_buffer(&self.material_buffer.inner_buf, 0, bytemuck::bytes_of(&uni));
-        debug!("Wrote material to buffer");
 
         MeshBufferStorage {
             buffer_vertex: vertex_buffer,
             buffer_indices,
-            material: Some(material),
+            material: Some(material.clone()),
         }
     }
 
@@ -511,13 +504,10 @@ impl Renderer for BasicRenderer {
         }
 
         if let Some(camera) = main_world.view_iter::<&mut CameraComponent>().next() {
-            let view_proj = self.inuse_camera.update_view_projection(&camera);
+            let uniform = self.inuse_camera.calc_view_projection(&camera);
-            let pos = camera.transform.translation;
+            //let pos = camera.transform.translation;
-            let uniform = CameraUniform {
+            //let uniform = CameraUniform::new(view_mat, *view_proj, pos);
-                view_proj: *view_proj,
+            self.camera_buffer.write_buffer(&self.queue, 0, &[uniform]);
-                view_pos: glam::Vec4::new(pos.x, pos.y, pos.z, 0.0),
-            };
-            self.queue.write_buffer(&self.camera_buffer, 0, bytemuck::cast_slice(&[uniform]));
         } else {
             warn!("Missing camera!");
         }
@@ -529,6 +519,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(&self.camera_buffer, &self.light_buffers, &self.depth_buffer_texture);
+
         let mut encoder = self.device.create_command_encoder(&wgpu::CommandEncoderDescriptor {
             label: Some("Basic Renderer's Encoder")
         });
@@ -588,10 +580,12 @@ 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, ]);
 
-                    render_pass.set_bind_group(2, &self.camera_bind_group, &[]);
+                    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(3, &self.light_buffers.bind_group_pair.bindgroup, &[]);
                     render_pass.set_bind_group(4, &self.material_buffer.bindgroup_pair.as_ref().unwrap().bindgroup, &[]);
         
+                    render_pass.set_bind_group(6, &self.light_cull_compute.light_indices_grid.bg_pair.bindgroup, &[]);
+
                     // if this mesh uses indices, use them to draw the mesh
                     if let Some((idx_type, indices)) = buffers.buffer_indices.as_ref() {
                         let indices_len = indices.count() as u32;
@@ -623,8 +617,15 @@ impl Renderer for BasicRenderer {
 
             // tell other things of updated resize
             self.surface.configure(&self.device, &self.config);
+            
+            let create_bindgroup = self.depth_buffer_texture.bindgroup_pair.is_some();
             self.depth_buffer_texture = RenderTexture::create_depth_texture(&self.device, &self.config, "Depth Buffer Texture");
+            if create_bindgroup {
+                self.depth_buffer_texture.create_bind_group(&self.device);
+            }
+            
             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,12 @@
 
 const max_light_count: u32 = 16u;
 
+const LIGHT_TY_DIRECTIONAL = 0u;
+const LIGHT_TY_POINT = 1u;
+const LIGHT_TY_SPOT = 2u;
+
+const ALPHA_CUTOFF = 0.1;
+
 struct VertexInput {
     @location(0) position: vec3<f32>,
     @location(1) tex_coords: vec2<f32>,
@@ -16,57 +22,33 @@ struct VertexOutput {
 }
 
 struct CameraUniform {
-    view_proj: mat4x4<f32>,
+    view: mat4x4<f32>,
-    view_pos: vec4<f32>,
+    inverse_projection: mat4x4<f32>,
-};
+    view_projection: mat4x4<f32>,
-
+    projection: mat4x4<f32>,
-struct PointLight {
-    position: vec4<f32>,
-    color: vec4<f32>,
-    
-    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>,
+    tile_debug: u32,
+};
+
+struct Light {
+    position: vec3<f32>,
+    light_ty: u32,
     direction: vec3<f32>,
+    enabled: u32,
     color: vec3<f32>,
     
-    cutoff: f32,
+    range: f32,
-    outer_cutoff: f32,
+    intensity: f32,
+    smoothness: 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>;
@@ -77,7 +59,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(
@@ -86,7 +68,7 @@ fn vs_main(
     var out: VertexOutput;
     
     out.tex_coords = model.tex_coords;
-    out.clip_position = u_camera.view_proj * u_model_transform * vec4<f32>(model.position, 1.0);
+    out.clip_position = u_camera.view_projection * u_model_transform * vec4<f32>(model.position, 1.0);
     
     // the normal mat is actually only a mat3x3, but there's a bug in wgpu: https://github.com/gfx-rs/wgpu-rs/issues/36
     let normal_mat = mat3x3(u_model_normal_matrix[0].xyz, u_model_normal_matrix[1].xyz, u_model_normal_matrix[2].xyz);
@@ -120,32 +102,96 @@ var t_specular: texture_2d<f32>;
 @group(5) @binding(1)
 var s_specular: sampler;
 
+@group(6) @binding(0)
+var<storage, read_write> u_light_indices: array<u32>;
+@group(6) @binding(1)
+var t_light_grid: texture_storage_2d<rg32uint, read_write>; // vec2<u32>
+
 @fragment
 fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
-    let object_color: vec4<f32> = textureSample(t_diffuse, s_diffuse, in.tex_coords);
+    /*let object_color: vec4<f32> = textureSample(t_diffuse, s_diffuse, in.tex_coords);
     let specular_color: vec3<f32> = textureSample(t_specular, s_specular, in.tex_coords).xyz;
 
-    // this needs to be 0.0 for the math
+    var light_res = vec3<f32>(0.0);
-    //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);
+    for (var i = 0u; i < u_lights.light_count; i++) {
+        var light = u_lights.data[i];
 
-    for (var i = 0u; i < u_lights.point_light_count; i++) {
+        if (light.light_ty == LIGHT_TY_DIRECTIONAL) {
-        light_res += blinn_phong_point_light(in.world_position, in.world_normal, u_lights.point_lights[i], u_material, specular_color);
+            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);
-    for (var i = 0u; i < u_lights.spot_light_count; i++) {
+        } else if (light.light_ty == LIGHT_TY_SPOT) {
-        light_res += blinn_phong_spot_light(in.world_position, in.world_normal, u_lights.spot_lights[i], u_material, specular_color);
+            light_res += blinn_phong_spot_light(in.world_position, in.world_normal, light, 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);*/
+    
+
+    if (u_camera.tile_debug == 1u) {
+        return debug_grid(in);
+    }
+
+    let object_color: vec4<f32> = textureSample(t_diffuse, s_diffuse, in.tex_coords);
+    let specular_color: vec3<f32> = textureSample(t_specular, s_specular, in.tex_coords).xyz;
+    var light_res = vec3<f32>(0.0);
+
+    if (object_color.a < ALPHA_CUTOFF) {
+        discard;
+    }
+
+    let tile_index = vec2<u32>(floor(in.clip_position.xy / 16.0));
+    let tile: vec2<u32> = textureLoad(t_light_grid, tile_index).xy;
+
+    let light_offset = tile.x;
+    let light_count = tile.y;
+
+    for (var i = 0u; i < light_count; i++) {
+        let light_index = u_light_indices[light_offset + i];
+        let light: Light = u_lights.data[light_index];
+
+        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);
+        }
+    }
+
+    let light_object_res = light_res * (object_color.xyz);
     return vec4<f32>(light_object_res, object_color.a);
+
+    //return debug_grid(in);
 }
 
-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 debug_grid(in: VertexOutput) -> vec4<f32> {
+    let tile_index_float: vec2<f32> = in.clip_position.xy / 16.0;
+    let tile_index = vec2<u32>(floor(tile_index_float));
+    let tile: vec2<u32> = textureLoad(t_light_grid, tile_index).xy;
+
+    // detect where the line grids would be at    
+    let x = tile_index_float.x - trunc(tile_index_float.x);
+    let y = tile_index_float.y - trunc(tile_index_float.y);
+    let ta: bool = x < 0.05 || y < 0.05;
+    let tb: bool = x > 0.95 || y > 0.95;
+
+    let ratio = f32(tile.y) / f32(u_lights.light_count);
+    return vec4<f32>(ratio, ratio, ratio, 1.0);
+
+    /* if ( ta || tb ) {
+        return vec4<f32>(0.0, 0.0, 0.0, 1.0);
+    } else {
+        return vec4<f32>(f32(tile_index.x) / 50.0, f32(tile_index.y) / 38.0, 0.0, 1.0);
+    } */
+}
+
+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;
+    let camera_view_pos = u_camera.position;
 
     //// Ambient light ////
     var ambient_color = light_color * material.ambient.xyz * material.diffuse.xyz;
@@ -165,17 +211,17 @@ 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;
+    return (ambient_color + diffuse_color + specular_color) * dir_light.intensity;
 }
 
-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;
+    let camera_view_pos = u_camera.position;
 
     //// Ambient light ////
     var ambient_color = light_color * material.ambient.xyz * material.diffuse.xyz;
@@ -195,67 +241,64 @@ 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 ////
     let distance = length(light_pos - world_pos);
-    let attenuation = 1.0 / (point_light.constant + point_light.linear * distance + 
+    let attenuation = 1.0 - smoothstep(point_light.range * point_light.smoothness, point_light.range, distance);
-    	point_light.quadratic * (distance * distance));
 
-    //// end of point light attenuation ////
+    ambient_color *= attenuation;
-    
+    diffuse_color *= attenuation;
-    ambient_color *= point_light.ambient * attenuation;
+    specular_color *= attenuation;
-    diffuse_color *= point_light.diffuse * attenuation;
-    specular_color *= point_light.specular * 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;
+    let light_color = spot_light.color;
-    let light_pos = spot_light.position.xyz;
+    let light_pos = spot_light.position;
-    let camera_view_pos = u_camera.view_pos.xyz;
+    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;
 
-    //if (theta > spot_light.cutoff) {
+    //// diffuse ////
-        var ambient_color = light_color * material.ambient.xyz * material.diffuse.xyz;
+    //let light_dir = normalize(light_pos - world_pos);
 
-        //// diffuse ////
+    let diffuse_strength = max(dot(world_norm, light_dir), 0.0);
-        //let light_dir = normalize(light_pos - world_pos);
+    var diffuse_color = light_color * (diffuse_strength * material.diffuse.xyz);
+    //// end of diffuse ////
 
-        let diffuse_strength = max(dot(world_norm, light_dir), 0.0);
+    //// specular ////
-        var diffuse_color = light_color * (diffuse_strength * material.diffuse.xyz);
+    let view_dir = normalize(camera_view_pos - world_pos);
-        //// end of diffuse ////
+    let half_dir = normalize(view_dir + light_dir);
 
-        //// specular ////
+    let specular_strength = pow(max(dot(world_norm, half_dir), 0.0), material.shininess);
-        let view_dir = normalize(camera_view_pos - world_pos);
+    var specular_color = specular_strength * (light_color * specular_factor);
-        let half_dir = normalize(view_dir + light_dir);
+    //// end of specular ////
     
-        let specular_strength = pow(max(dot(world_norm, half_dir), 0.0), material.shininess);
+    //// spot light soft edges ////
-        var specular_color = specular_strength * (light_color * specular_factor);
+    let min_cos = cos(spot_light.spot_cutoff);
-        //// end of specular ////
+    let max_cos = lerp(min_cos, 1.0, 0.5);
+    let cos_angle = dot(spot_light.direction, -light_dir);
+    let cone = smoothstep(min_cos, max_cos, cos_angle);
+    //// end of spot light soft edges ////
 
-        //// spot light soft edges ////
+    //// spot light attenuation ////
-        let theta = dot(light_dir, normalize(-spot_light.direction));
+    let distance = length(light_pos - world_pos);
-        let epsilon = spot_light.cutoff - spot_light.outer_cutoff;
+    let attenuation = calc_attenuation(spot_light, distance);
-        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 ////
+    ambient_color *= attenuation * spot_light.intensity * cone;
-        let distance = length(light_pos - world_pos);
+    diffuse_color *= attenuation * spot_light.intensity * cone;
-        let attenuation = 1.0 / (spot_light.constant + spot_light.linear * distance + 
+    specular_color *= attenuation * spot_light.intensity * cone;
-            spot_light.quadratic * (distance * distance));
+    //// end of spot light attenuation ////
-
-        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;
+    return /*ambient_color +*/ diffuse_color + specular_color;
-    /*} else {
+}
-        return vec3<f32>(0.0);
+
-    }*/
+fn calc_attenuation(light: Light, distance: f32) -> f32 {
+    return 1.0 - smoothstep(light.range * light.smoothness, light.range, distance);
+}
+
+fn lerp(start: f32, end: f32, alpha: f32) -> f32 {
+    return (start + (end - start) * alpha);
 }

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

@@ -0,0 +1,329 @@
+const BLOCK_SIZE: u32 = 16u;
+const MAX_TILE_VISIBLE_LIGHTS: u32 = 1024u;
+
+const LIGHT_TY_DIRECTIONAL = 0u;
+const LIGHT_TY_POINT = 1u;
+const LIGHT_TY_SPOT = 2u;
+
+type vec2f = vec2<f32>;
+type vec3f = vec3<f32>;
+type vec4f = vec4<f32>;
+
+struct CameraUniform {
+    view: mat4x4<f32>,
+    inverse_projection: mat4x4<f32>,
+    view_projection: mat4x4<f32>,
+    projection: mat4x4<f32>,
+    position: vec3f,
+    tile_debug: u32,
+};
+
+struct Light {
+    position: vec3f,
+    light_ty: u32,
+    direction: vec3f,
+    enabled: u32,
+    color: vec3f,
+    
+    range: f32,
+    intensity: f32,
+    smoothness: f32,
+
+    spot_cutoff: f32,
+    spot_outer_cutoff: f32,
+};
+
+struct Lights {
+    light_count: u32,
+    data: array<Light>,
+};
+
+struct Cone {
+    tip: vec3f,
+    height: f32,
+    direction: vec3f,
+    radius: f32,
+}
+
+struct Plane {
+    normal: vec3f,
+    origin_distance: f32,
+}
+
+var<workgroup> wg_min_depth: atomic<u32>;
+var<workgroup> wg_max_depth: atomic<u32>;
+var<workgroup> wg_light_index_start: atomic<u32>;
+var<workgroup> wg_frustum_planes: array<Plane, 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>;
+
+@group(0) @binding(0)
+var t_depthmap: texture_depth_2d;
+@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 t_light_grid: texture_storage_2d<rg32uint, read_write>;
+@group(3) @binding(2)
+var<storage, read_write> u_light_index_counter: atomic<u32>;
+
+@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,
+) {
+    // 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<u32>(global_invocation_id.xy);
+    var depth_float: f32 = textureLoad(t_depthmap, tex_coord, 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);
+
+    workgroupBarrier();
+
+    // convert them back into floats
+    var min_depth: f32 = bitcast<f32>(wg_min_depth);
+    var max_depth: f32 = bitcast<f32>(wg_max_depth);
+    
+    // Create the frustum planes that will be used for this time
+    if (local_invocation_index == 0u) {
+        // this algorithm is adapted from Google's filament:
+        // https://github.com/google/filament/blob/3644e7f80827f1cd2caef4a21e410a2243eb6e84/filament/src/Froxelizer.cpp#L402C57-L402C73
+        let tile_width_clip_space = f32(2u * BLOCK_SIZE) / f32(u_screen_size.x);
+        let tile_height_clip_space = f32(2u * BLOCK_SIZE) / f32(u_screen_size.y);
+
+        let tr_projection = transpose(u_camera.projection);
+
+        var planes: array<vec4f, 4>;
+
+        // left plane
+        {
+            let x = (f32(workgroup_id.x) * tile_width_clip_space) - 1.0;
+            let p = tr_projection * vec4f(-1.0, 0.0, 0.0, x);
+            planes[0] = -vec4f(normalize(p.xyz), 0.0);
+        }
+
+        // right plane
+        {
+            let x = (f32(workgroup_id.x + 1u) * tile_width_clip_space) - 1.0;
+            let p = tr_projection * vec4f(-1.0, 0.0, 0.0, x);
+            planes[1] = vec4f(normalize(p.xyz), 0.0);
+        }
+
+        // top plane
+        {
+            let y = (f32(workgroup_id.y) * tile_height_clip_space) - 1.0;
+            let p = tr_projection * vec4f(0.0, 1.0, 0.0, y);
+            planes[2] = -vec4f(normalize(p.xyz), 0.0);
+        }
+
+        // bottom plane
+        {
+            let y = (f32(workgroup_id.y + 1u) * tile_height_clip_space) - 1.0;
+            let p = tr_projection * vec4f(0.0, 1.0, 0.0, y);
+            planes[3] = vec4f(normalize(p.xyz), 0.0);
+        }
+
+        wg_frustum_planes[0] = Plane(planes[0].xyz, planes[0].w);
+        wg_frustum_planes[1] = Plane(planes[1].xyz, planes[1].w);
+        wg_frustum_planes[2] = Plane(planes[2].xyz, planes[2].w);
+        wg_frustum_planes[3] = Plane(planes[3].xyz, planes[3].w);
+
+        wg_frustum_planes[4] = Plane(vec3f(0.0, 0.0, -1.0), -min_depth);
+        wg_frustum_planes[5] = Plane(vec3f(0.0, 0.0, 1.0), -max_depth);
+    }
+
+    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.
+    for (var i = local_invocation_index; i < u_lights.light_count; i += BLOCK_SIZE * BLOCK_SIZE) {
+        let light_index = i;
+
+        let light = u_lights.data[light_index];
+
+        if (light.enabled == 1u) {
+            let position_vs = (u_camera.view * vec4f(light.position, 1.0)).xyz;
+
+            if (light.light_ty == LIGHT_TY_DIRECTIONAL) {
+                add_light(light_index);
+            } else if (light.light_ty == LIGHT_TY_POINT
+                    && sphere_inside_frustrum(wg_frustum_planes, position_vs, light.range)) {
+                // TODO: add the light to the transparent geometry list
+
+                if (!sphere_inside_plane(position_vs, light.range, wg_frustum_planes[4])) {
+                    add_light(light_index);
+                }
+            } else if (light.light_ty == LIGHT_TY_SPOT) {
+                let dir_vs = (u_camera.view * vec4f(light.direction, 1.0)).xyz;
+                let cone_radius = tan(light.spot_cutoff) * light.range;
+                let cone = Cone(position_vs, light.range, dir_vs, cone_radius);
+
+                if (cone_inside_frustum(cone, wg_frustum_planes)) {
+                    // TODO: add the light to the transparent geometry list
+
+                    add_light(light_index);
+                    if (!cone_inside_plane(cone, wg_frustum_planes[4])) {
+                    }
+                }
+            }
+        }
+    }
+
+    workgroupBarrier();
+
+    // Update the global memory with the visible light buffer.
+
+    // first update the light grid on the first thread
+    if (local_invocation_index == 0u) {
+        wg_light_index_start = atomicAdd(&u_light_index_counter, wg_visible_light_count);
+        textureStore(t_light_grid, workgroup_id.xy, vec4<u32>(wg_light_index_start, wg_visible_light_count, 0u, 1u));
+
+        // TODO: store light grid for transparent geometry
+    }
+
+    workgroupBarrier();
+
+    // now update the light index list on all threads.
+    for (var i = local_invocation_index; i < wg_visible_light_count; i += BLOCK_SIZE * BLOCK_SIZE) {
+        u_light_indices[wg_light_index_start + i] = wg_visible_light_indices[i];
+    }
+}
+
+/// Add a light to the visible light indicies list.
+/// Returns a boolean indicating if the light was added.
+fn add_light(light_index: u32) -> bool {
+    //var offset: u32 = wg_visible_light_count;
+
+    if (wg_visible_light_count < MAX_TILE_VISIBLE_LIGHTS) {
+        let offset = atomicAdd(&wg_visible_light_count, 1u);
+        wg_visible_light_indices[offset] = light_index;
+        return true;
+    }
+
+    return false;
+}
+
+fn sphere_inside_frustrum(frustum: array<Plane, 6>, sphere_origin: vec3f, 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: vec3f, radius: f32, plane: Plane) -> bool {
+    return dot(plane.normal, sphere_origin) - plane.origin_distance < -radius;
+}
+
+fn clip_to_view(clip: vec4f) -> vec4f {
+    // view space position
+    var view = u_camera.inverse_projection * clip;
+
+    // perspective projection
+    return view / view.w;
+}
+
+fn screen_to_view(screen: vec4f) -> vec4f {
+    // convert to normalized texture coordinates
+    let tex_coord = screen.xy / vec2<f32>(u_screen_size);
+
+    // convert to clip space
+    let clip = vec4f( vec2<f32>(tex_coord.x, 1.0 - tex_coord.y) * 2.0 - 1.0, screen.z, screen.w);
+    
+    return clip_to_view(clip);
+}
+
+/// Compute a plane from 3 noncollinear points that form a triangle.
+/// This equation assumes a right-handed (counter-clockwise winding order) 
+/// coordinate system to determine the direction of the plane normal.
+fn compute_plane(p0: vec3f, p1: vec3f, p2: vec3f) -> Plane {
+    let v0 = p1 - p0;
+    let v2 = p2 - p0;
+
+    let normal = vec4f(normalize(cross(v0, v2)), 0.0);
+
+    // find the distance to the origin
+    let distance = dot(normal.xyz, p0);
+
+    return Plane(normal.xyz, distance);
+}
+
+fn point_inside_plane(point: vec3f, plane: Plane) -> bool {
+    return dot(plane.normal, point) + plane.origin_distance < 0.0;
+}
+
+fn point_intersect_plane(point: vec3f, plane: Plane) -> f32 {
+    return dot(plane.normal, point) + plane.origin_distance;
+}
+
+/// 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: Plane) -> bool {
+    let dir = cone.direction;
+    let furthest_direction = cross(cross(plane.normal, dir), dir);
+    let furthest = cone.tip + dir * cone.height - furthest_direction * 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(furthest, plane);
+}
+
+fn cone_inside_frustum(cone: Cone, frustum: array<Plane, 6>) -> bool {
+    var frustum = frustum;
+    for (var i = 0u; i < 4u; i++) {
+        // TODO: better cone checking
+        if (sphere_inside_plane(cone.tip, cone.radius, frustum[i])) {
+            return false;
+        }
+    }
+
+    return true;
+}

lyra-game/src/render/texture.rs

@@ -1,4 +1,4 @@
-use std::sync::Arc;
+use std::rc::Rc;
 
 use image::GenericImageView;
 use lyra_resource::{FilterMode, ResHandle, Texture, WrappingMode};
@@ -44,7 +44,7 @@ impl RenderTexture {
         })
     }
 
-    fn create_bind_group_pair(device: &wgpu::Device, layout: Arc<wgpu::BindGroupLayout>, view: &wgpu::TextureView, sampler: &wgpu::Sampler) -> BindGroupPair {
+    fn create_bind_group_pair(device: &wgpu::Device, layout: Rc<wgpu::BindGroupLayout>, view: &wgpu::TextureView, sampler: &wgpu::Sampler) -> BindGroupPair {
         let bg = device.create_bind_group(
             &wgpu::BindGroupDescriptor {
                 layout: &layout,
@@ -68,12 +68,12 @@ impl RenderTexture {
         }
     }
 
-    pub fn from_bytes(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Arc<wgpu::BindGroupLayout>, bytes: &[u8], label: &str) -> anyhow::Result<Self> {
+    pub fn from_bytes(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Rc<wgpu::BindGroupLayout>, bytes: &[u8], label: &str) -> anyhow::Result<Self> {
         let img = image::load_from_memory(bytes)?;
         Self::from_image(device, queue, bg_layout, &img, Some(label))
     }
 
-    pub fn from_image(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Arc<wgpu::BindGroupLayout>, img: &image::DynamicImage, label: Option<&str>) -> anyhow::Result<Self> {
+    pub fn from_image(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Rc<wgpu::BindGroupLayout>, img: &image::DynamicImage, label: Option<&str>) -> anyhow::Result<Self> {
         let rgba = img.to_rgba8();
         let dimensions = img.dimensions();
 
@@ -134,7 +134,7 @@ impl RenderTexture {
         })
     }
 
-    pub fn from_resource(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Arc<wgpu::BindGroupLayout>, texture_res: &ResHandle<Texture>, label: Option<&str>) -> anyhow::Result<Self> {
+    pub fn from_resource(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Rc<wgpu::BindGroupLayout>, texture_res: &ResHandle<Texture>, label: Option<&str>) -> anyhow::Result<Self> {
         let texture_ref = texture_res.data_ref().unwrap();
         let img = texture_ref.image.data_ref().unwrap();
 
@@ -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:

lyra-game/src/scene/camera.rs

@@ -7,6 +7,7 @@ pub struct CameraComponent {
     pub transform: Transform,
     pub fov: Angle,
     pub mode: CameraProjectionMode,
+    pub debug: bool,
 }
 
 impl Default for CameraComponent {
@@ -15,6 +16,7 @@ impl Default for CameraComponent {
             transform: Transform::default(),
             fov: Angle::Degrees(45.0),
             mode: CameraProjectionMode::Perspective,
+            debug: false,
         }
     }
 }

lyra-game/src/scene/free_fly_camera.rs

@@ -1,28 +1,103 @@
-use lyra_ecs::Component;
+use glam::{EulerRot, Quat, Vec3};
+use lyra_ecs::{query::{Res, View}, Component};
 
-use crate::{math::{Angle, Transform}, render::camera::CameraProjectionMode};
+use crate::{game::Game, input::ActionHandler, plugin::Plugin, DeltaTime};
+
+use super::CameraComponent;
+
+pub const ACTLBL_MOVE_UP_DOWN: &str = "MoveUpDown";
+pub const ACTLBL_MOVE_LEFT_RIGHT: &str = "MoveLeftRight";
+pub const ACTLBL_MOVE_FORWARD_BACKWARD: &str = "MoveForwardBackward";
+pub const ACTLBL_LOOK_LEFT_RIGHT: &str = "LookLeftRight";
+pub const ACTLBL_LOOK_UP_DOWN: &str = "LookUpDown";
+pub const ACTLBL_LOOK_ROLL: &str = "LookRoll";
 
 #[derive(Clone, Component)]
 pub struct FreeFlyCamera {
-    pub transform: Transform,
-    pub fov: Angle,
-    pub mode: CameraProjectionMode,
     pub speed: f32,
+    pub slow_speed_factor: f32,
+    pub look_speed: f32,
+    pub mouse_sensitivity: f32,
+    pub look_with_keys: bool,
 }
 
 impl Default for FreeFlyCamera {
     fn default() -> Self {
-        Self::new()
+        Self {
+            speed: 4.0,
+            slow_speed_factor: 0.25,
+            look_speed: 0.5,
+            mouse_sensitivity: 0.9,
+            look_with_keys: false,
+        }
     }
 }
 
 impl FreeFlyCamera {
-    pub fn new() -> Self {
+    #[allow(dead_code)]
+    pub fn new(speed: f32, slow_speed_factor: f32, look_speed: f32, mouse_sensitivity: f32, look_with_keys: bool) -> Self {
         Self {
-            transform: Transform::default(),
+            speed,
-            fov: Angle::Degrees(45.0),
+            slow_speed_factor,
-            mode: CameraProjectionMode::Perspective,
+            look_speed,
-            speed: 1.5,
+            mouse_sensitivity,
+            look_with_keys,
         }
     }
 }
+
+pub fn free_fly_camera_controller(delta_time: Res<DeltaTime>, handler: Res<ActionHandler>, view: View<(&mut CameraComponent, &FreeFlyCamera)>) -> anyhow::Result<()> {
+    let delta_time = **delta_time;
+    for (mut cam, fly) in view.into_iter() {
+        let forward = cam.transform.forward();
+        let left = cam.transform.left();
+        let up = Vec3::Y;
+
+        let move_y = handler.get_axis_modifier(ACTLBL_MOVE_UP_DOWN).unwrap_or(0.0);
+        let move_x = handler.get_axis_modifier(ACTLBL_MOVE_LEFT_RIGHT).unwrap_or(0.0);
+        let move_z = handler.get_axis_modifier(ACTLBL_MOVE_FORWARD_BACKWARD).unwrap_or(0.0);
+
+        let mut velocity = Vec3::ZERO;
+        velocity += move_y * up;
+        velocity += move_x * left;
+        velocity += move_z * forward;
+
+        if velocity != Vec3::ZERO {
+            cam.transform.translation += velocity.normalize() * fly.speed * delta_time; // TODO: speeding up
+        }
+
+        let motion_x = handler.get_axis_modifier(ACTLBL_LOOK_LEFT_RIGHT).unwrap_or(0.0);
+        let motion_y = handler.get_axis_modifier(ACTLBL_LOOK_UP_DOWN).unwrap_or(0.0);
+        let motion_z = handler.get_axis_modifier(ACTLBL_LOOK_ROLL).unwrap_or(0.0);
+        
+        let mut camera_rot = Vec3::ZERO;
+        camera_rot.y -= motion_x * fly.mouse_sensitivity;
+        camera_rot.x -= motion_y * fly.mouse_sensitivity;
+        camera_rot.z -= motion_z * fly.mouse_sensitivity;
+
+        if camera_rot != Vec3::ZERO {
+            let look_velocity = camera_rot * fly.look_speed * delta_time;
+
+            let (mut y, mut x, _) = cam.transform.rotation.to_euler(EulerRot::YXZ);
+            x += look_velocity.x;
+            y += look_velocity.y;
+            x = x.clamp(-1.54, 1.54);
+
+            // rotation is not commutative, keep this order to avoid unintended roll
+            cam.transform.rotation = Quat::from_axis_angle(Vec3::Y, y)
+                * Quat::from_axis_angle(Vec3::X, x);
+        }
+    }
+
+    Ok(())
+}
+
+/// A plugin that adds the free fly camera controller system to the world. It is expected that
+/// there is a [`FreeFlyCamera`] in the world, if there isn't, the camera would not move.
+pub struct FreeFlyCameraPlugin;
+
+impl Plugin for FreeFlyCameraPlugin {
+    fn setup(&self, game: &mut Game) {
+        game.with_system("free_fly_camera_system", free_fly_camera_controller, &[]);
+    }
+}

rust-toolchain.toml

@@ -1,5 +1,4 @@
 [toolchain]
-channel = "nightly-2023-11-21"
+channel = "nightly"
-#components = [ "rustfmt", "rustc-dev" ]
+date = "2023-11-21"
 targets = [ "x86_64-unknown-linux-gnu" ]
-#profile = "minimal"

shell.nix

@@ -1,24 +1,33 @@
-{ pkgs ? import <nixpkgs> { } }:
+let
+  moz_overlay = import (builtins.fetchTarball https://github.com/mozilla/nixpkgs-mozilla/archive/master.tar.gz);
+  nixpkgs = import <nixpkgs> { overlays = [ moz_overlay ]; };
+in
+  with nixpkgs;
+  stdenv.mkDerivation rec {
+    name = "lyra_engine_dev";
+    nativeBuildInputs = [
+      pkg-config
+      openssl
+      wasm-pack
+      trunk
+      valgrind
+      heaptrack
+      mold
+      udev
+      lua5_4_compat
+      ((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
       
-with pkgs;
+    ];
-
+    LD_LIBRARY_PATH = lib.makeLibraryPath buildInputs;
-mkShell rec {
+  }
-  nativeBuildInputs = [
-    pkg-config
-    openssl
-    wasm-pack
-    trunk
-    valgrind
-    heaptrack
-    mold
-    udev
-    lua5_4_compat
-  ];
-  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;
-}