render: get some lights showing up with tiled forward rendering

For some reason there's weird square in the light source, and the dynamic light is only applied to the top left tile

examples/testbed/src/free_fly_camera.rs

@@ -25,8 +25,8 @@ impl Default for FreeFlyCamera {
         Self {
             speed: 4.0,
             slow_speed_factor: 0.25,
-            look_speed: 0.3,
-            mouse_sensitivity: 1.0,
+            look_speed: 0.5,
+            mouse_sensitivity: 0.9,
             look_with_keys: false,
         }
     }

examples/testbed/src/main.rs

@@ -1,6 +1,6 @@
 use std::{cell::Ref, ptr::NonNull};
 
-use lyra_engine::{assets::gltf::Gltf, ecs::{query::{QueryBorrow, 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, SpotLight}, scene::CameraComponent, DeltaTime};
+use lyra_engine::{assets::gltf::Gltf, change_tracker::Ct, ecs::{query::{QueryBorrow, 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, DeltaTime};
 use lyra_engine::assets::ResourceManager;
 
 mod free_fly_camera;
@@ -93,15 +93,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 crate_gltf = resman.request::<Gltf>("assets/crate/crate.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 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()
@@ -151,29 +152,33 @@ async fn main() {
             ));
         }
 
-        /* {
-            let mut light_tran = Transform::from_xyz(-3.5, 0.2, -4.5);
-            light_tran.scale = Vec3::new(0.5, 0.5, 0.5);
+        {
+            //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),
+                PointLight {
+                    color: Vec3::new(0.0, 0.0, 1.0),
+
+                    intensity: 3.3,
 
                     constant: 1.0,
-                    linear: 0.007,
-                    quadratic: 0.0002,
+                    linear: 0.09,
+                    quadratic: 0.032,
                     
-                    ambient: 0.0,
-                    diffuse: 7.0,
-                    specular: 1.0,
+                    ambient: 0.2,
+                    diffuse: 1.0,
+                    specular: 1.3,
                 },
-                Transform::from(light_tran),
+                Transform::new(
+                    Vec3::new(-5.0, 1.0, -1.28), 
+                    Quat::IDENTITY,
+                    Vec3::new(0.25, 0.25, 0.25),
+                ),
                 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((

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/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::mem;
+use std::{collections::{HashMap, VecDeque}, marker::PhantomData};
 
 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,8 +101,7 @@ impl<U: Default + bytemuck::Pod + bytemuck::Zeroable> LightBuffer<U> {
 
 pub(crate) struct LightUniformBuffers {
     pub buffer: wgpu::Buffer,
-    pub bindgroup_layout: wgpu::BindGroupLayout,
-    pub bindgroup: wgpu::BindGroup,
+    pub bind_group_pair: BindGroupPair,
     pub light_indexes: HashMap<Entity, u32>,
     pub current_light_idx: u32,
 }
@@ -159,8 +158,7 @@ impl LightUniformBuffers {
 
         Self {
             buffer,
-            bindgroup_layout,
-            bindgroup,
+            bind_group_pair: BindGroupPair::new(bindgroup, bindgroup_layout),
             light_indexes: Default::default(),
             current_light_idx: 0,
         }

lyra-game/src/render/light_cull_compute.rs

@@ -1,47 +1,36 @@
-use std::{borrow::Cow, mem, num::NonZeroU32, ptr::NonNull, rc::Rc};
+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::BufferWrapper, renderer::RenderPass, texture::RenderTexture};
+use super::{light::LightUniformBuffers, render_buffer::{BindGroupPair, BufferWrapper}, texture::RenderTexture};
 
-struct LightIndicesGridBuffer {
+pub(crate) struct LightIndicesGridBuffer {
     indices_buffer: wgpu::Buffer,
     grid_texture: wgpu::Texture,
     grid_texture_view: wgpu::TextureView,
-    bind_group_layout: wgpu::BindGroupLayout,
-    bind_group: wgpu::BindGroup,
+    pub bg_pair: BindGroupPair,
 }
 
 pub(crate) struct LightCullCompute {
     device: Rc<wgpu::Device>,
     queue: Rc<wgpu::Queue>,
     pipeline: ComputePipeline,
-    lights: NonNull<LightUniformBuffers>,
-    camera: NonNull<BufferWrapper>,
-    light_indices_grid: LightIndicesGridBuffer,
+    pub light_indices_grid: LightIndicesGridBuffer,
     screen_size_buffer: BufferWrapper,
-    depth_tex: NonNull<RenderTexture>,
+    workgroup_size: glam::UVec2,
 }
 
 impl LightCullCompute {
-    fn create_grid(device: &wgpu::Device, screen_size: PhysicalSize<u32>) -> LightIndicesGridBuffer {
-        let limits = device.limits();
-        //let max_buffer_sizes = (limits.max_uniform_buffer_binding_size as u64) / 2;
-
-        /* let light_indices_buffer = device.create_buffer(
-            &wgpu::BufferDescriptor {
-                label: Some("B_LightIndices"),
-                usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
-                size: (mem::size_of::<u32>() * 16 * 16) as u64,
-                mapped_at_creation: false,
-            }
-        ); */
+    fn create_grid(device: &wgpu::Device, screen_size: PhysicalSize<u32>, workgroup_size: glam::UVec2) -> LightIndicesGridBuffer {
+        let mut contents = Vec::<u8>::new();
+        let contents_len = workgroup_size.x * workgroup_size.y * mem::size_of::<u8>() as u32;
+        contents.resize(contents_len as _, 0);
 
         let light_indices_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
             label: Some("B_LightIndices"),
-            contents: &[0; mem::size_of::<u32>() * 16 * 16],
+            contents: &contents,
             usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
         });
 
@@ -128,12 +117,11 @@ impl LightCullCompute {
             indices_buffer: light_indices_buffer,
             grid_texture,
             grid_texture_view,
-            bind_group_layout: light_indices_bg_layout,
-            bind_group: light_indices_bg,
+            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: &LightUniformBuffers, camera_buffers: &BufferWrapper, depth_texture: &mut RenderTexture) -> Self {
+    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")
@@ -148,18 +136,18 @@ impl LightCullCompute {
             source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(shader_src)),
         });
 
-        let light_grid = Self::create_grid(&device, screen_size);
+        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, screen_size, 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_texture.bindgroup_pair.as_ref().unwrap().layout,
                 &depth_tex_pair.layout,
-                camera_buffers.bindgroup_layout().unwrap(),
-                &lights.bindgroup_layout,
-                &light_grid.bind_group_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: &[],
@@ -176,22 +164,19 @@ impl LightCullCompute {
             device,
             queue,
             pipeline,
-            lights: NonNull::from(lights),
-            camera: NonNull::from(camera_buffers),
             light_indices_grid: light_grid,
             screen_size_buffer,
-            depth_tex: NonNull::from(depth_texture),
+            workgroup_size,
         }
     }
 
-    pub fn update_screen_size(&self, size: PhysicalSize<u32>) {
+    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);
     }
 
-    pub fn compute(&mut self) {
-        //self.queue.write_buffer(&self.light_indices_grid.indices_buffer, 0, &[0; mem::size_of::<u32>() * 16 * 16]);
-
+    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"),
         });
@@ -203,19 +188,13 @@ impl LightCullCompute {
 
             pass.set_pipeline(&self.pipeline);
 
-            let depth = unsafe { self.depth_tex.as_ref() };
-            pass.set_bind_group(0, depth.bind_group(), &[]);
-
-            let cam = unsafe { self.camera.as_ref() };
-            pass.set_bind_group(1, cam.bindgroup(), &[]);
-
-            let lights = unsafe { self.lights.as_ref() };
-            pass.set_bind_group(2, &lights.bindgroup, &[]);
-
-            pass.set_bind_group(3, &self.light_indices_grid.bind_group, &[]);
+            pass.set_bind_group(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(16, 16, 1);
+            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);

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/render_buffer.rs

@@ -1,4 +1,4 @@
-use std::{num::NonZeroU32, ops::Deref, sync::Arc};
+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,
         });
@@ -43,7 +43,7 @@ impl BindGroupPair {
     pub fn new(bindgroup: wgpu::BindGroup, layout: wgpu::BindGroupLayout) -> Self {
         Self {
             bindgroup,
-            layout: Arc::new(layout),
+            layout: Rc::new(layout),
         }
     }
 }
@@ -92,7 +92,7 @@ impl BufferWrapper {
     /// 
     /// 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.deref())
+        self.bindgroup_pair.as_ref().map(|bg| &*bg.layout)
     }
 
     /// Queue's the data to be written to `buffer` starting at `offset`.
@@ -272,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 {
@@ -288,7 +288,7 @@ impl BufferWrapperBuilder {
 
                 BindGroupPair {
                     bindgroup: bg,
-                    layout: bg_layout,
+                    layout: Rc::new(bg_layout),
                 }
             }
         };

lyra-game/src/render/renderer.rs

@@ -59,7 +59,7 @@ struct MeshBufferStorage {
 
     //#[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)
@@ -88,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>,
 
@@ -99,7 +100,7 @@ pub struct BasicRenderer {
     camera_buffer: BufferWrapper,
     //camera_bind_group: wgpu::BindGroup,
 
-    bgl_texture: Arc<BindGroupLayout>,
+    bgl_texture: Rc<BindGroupLayout>,
     default_texture: RenderTexture,
     depth_buffer_texture: RenderTexture,
 
@@ -172,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 {
@@ -223,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,
@@ -247,8 +249,10 @@ impl BasicRenderer {
             vec![super::vertex::Vertex::desc(),],
             vec![&s.bgl_texture, &s.transform_buffers.bindgroup_layout, 
             s.camera_buffer.bindgroup_layout().unwrap(),
-            &s.light_buffers.bindgroup_layout, &s.material_buffer.bindgroup_pair.as_ref().unwrap().layout,
-            &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
@@ -347,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")
-            .data_ref()
+            .expect("Material resource not loaded yet");
+        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()),
         }
     }
 
@@ -509,7 +519,7 @@ 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.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")
@@ -571,9 +581,11 @@ impl Renderer for BasicRenderer {
                     render_pass.set_bind_group(1, bindgroup, &[ offset, offset, ]);
 
                     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;
@@ -605,7 +617,13 @@ 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,9 +2,9 @@
 
 const max_light_count: u32 = 16u;
 
-const light_ty_directional = 0u;
-const light_ty_point = 1u;
-const light_ty_spot = 2u;
+const LIGHT_TY_DIRECTIONAL = 0u;
+const LIGHT_TY_POINT = 1u;
+const LIGHT_TY_SPOT = 2u;
 
 struct VertexInput {
     @location(0) position: vec3<f32>,
@@ -97,41 +97,79 @@ 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
-    //u_lights.directional_light.direction.w = 0.0;
-
     var light_res = vec3<f32>(0.0);
 
     for (var i = 0u; i < u_lights.light_count; i++) {
         var light = u_lights.data[i];
 
-        if (light.light_ty == light_ty_directional) {
+        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) {
+        } 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) {
+        } else if (light.light_ty == LIGHT_TY_SPOT) {
             light_res += blinn_phong_spot_light(in.world_position, in.world_normal, light, u_material, specular_color);
         }
     }
 
-    /*var light_res = blinn_phong_dir_light(in.world_position, in.world_normal, u_lights.directional_light, u_material, specular_color);
-
-    for (var i = 0u; i < u_lights.point_light_count; i++) {
-        light_res += blinn_phong_point_light(in.world_position, in.world_normal, u_lights.point_lights[i], u_material, specular_color);
-    }
-
-    for (var i = 0u; i < u_lights.spot_light_count; i++) {
-        light_res += blinn_phong_spot_light(in.world_position, in.world_normal, u_lights.spot_lights[i], u_material, specular_color);
-    }*/
-
     let light_object_res = light_res * (object_color.xyz/* * u_material.diffuse.xyz*/);
 
+    return vec4<f32>(light_object_res, object_color.a);*/
+    
+    let tile_index = vec2<u32>(floor(in.clip_position.xy / 16.0));
+    let tile: vec2<u32> = textureLoad(t_light_grid, tile_index).xy;
+
+    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);
+
+    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 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;
+
+    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> {

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

@@ -1,9 +1,9 @@
-const block_size: i32 = 16;
-const max_tile_visible_lights: u32 = 1024u;
+const BLOCK_SIZE: i32 = 16;
+const MAX_TILE_VISIBLE_LIGHTS: u32 = 1024u;
 
-const light_ty_directional = 0u;
-const light_ty_point = 1u;
-const light_ty_spot = 2u;
+const LIGHT_TY_DIRECTIONAL = 0u;
+const LIGHT_TY_POINT = 1u;
+const LIGHT_TY_SPOT = 2u;
 
 // Possible computer shader inputs:
 //
@@ -44,13 +44,13 @@ var<workgroup> wg_max_depth: atomic<u32>;
 var<workgroup> wg_frustum_planes: array<vec4<f32>, 6>;
 
 // index list of visible light sources for this tile
-var<workgroup> wg_visible_light_indices: array<u32, max_tile_visible_lights>;
+var<workgroup> wg_visible_light_indices: array<u32, MAX_TILE_VISIBLE_LIGHTS>;
 var<workgroup> wg_visible_light_count: atomic<u32>;
 
 //var<workgroup> view_projection: mat4x4;
 
 @group(0) @binding(0)
-var t_depthmap: texture_2d<f32>;
+var t_depthmap: texture_depth_2d;
 @group(0) @binding(1)
 var s_depthmap: sampler;
 
@@ -62,11 +62,8 @@ var<storage, read> u_lights: Lights;
 
 @group(3) @binding(0)
 var<storage, read_write> u_light_indices: array<u32>;
-/*@group(3) @binding(1)
-var<uniform> u_light_grid: array<array<vec2<u32>>>;*/
-
 @group(3) @binding(1)
-var t_light_grid: texture_storage_2d<rg32uint, read_write>; // rg32uint = vec2<u32> or vec4<u32>(r, g, 0.0, 1.0) 
+var t_light_grid: texture_storage_2d<rg32uint, read_write>;
 
 @group(4) @binding(0)
 var<uniform> u_screen_size: vec2<u32>;
@@ -96,9 +93,8 @@ fn cs_main(
     workgroupBarrier();
 
     // step 1: calculate the minimum and maximum depth values for this tile (using the depth map)
-    var tex_coord = vec2<f32>(global_invocation_id.xy);
-    //var depth_float: f32 = textureSample(t_depthmap, s_depthmap, tex_coord).r;
-    var depth_float = 0.0;
+    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);
 
@@ -146,7 +142,7 @@ fn cs_main(
     // Process the lights detecting which ones to cull for this tile.
     // Processes 256 lights simultaniously, each on a thread in the workgroup. Requires multiple
     // iterations for more lights.
-    var thread_count = u32(block_size * block_size);
+    var thread_count = u32(BLOCK_SIZE * BLOCK_SIZE);
     var pass_count = (u_lights.light_count + thread_count - 1u) / thread_count;
     for (var i = 0u; i < pass_count; i++) {
         // find the light index to check on this thread, make sure we're not trying to test
@@ -160,18 +156,21 @@ fn cs_main(
         var position = light.position;
         var radius = light.range;
 
-        if (light.light_ty != light_ty_directional
+        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, radius)) {
             // TODO: add the light to the transparent geometry list
 
             // TODO: spotlights
             if (!sphere_inside_plane(position, radius, wg_frustum_planes[4])) {
-                var offset: u32 = wg_visible_light_count;
+                add_light(light_index);
+                /*var offset: u32 = wg_visible_light_count;
 
-                if (offset < max_tile_visible_lights) {
+                if (offset < MAX_TILE_VISIBLE_LIGHTS) {
                     atomicAdd(&wg_visible_light_count, 1u);
                     wg_visible_light_indices[offset] = light_index;
-                }
+                }*/
             }
         }
     }
@@ -182,28 +181,14 @@ fn cs_main(
 
     // first update the light grid on the first thread
     if (local_invocation_index == 0u) {
-        var offset = u32(index) * max_tile_visible_lights; // index in the global light list
-        //t_light_grid[workgroup_id.x][workgroup_id.y] = vec2<f32>(offset, wg_visible_light_count);
+        var offset = u32(index) * MAX_TILE_VISIBLE_LIGHTS; // index in the global light list
         textureStore(t_light_grid, workgroup_id.xy, vec4<u32>(offset, wg_visible_light_count, 0u, 1u));
-        
-        // TODO: update transparent light grid
-
-        /*var offset = index * max_tile_visible_lights; // position in the global light buffer
-        // update the light
-        for (var i = 0u; i < wg_visible_light_count; i++) {
-            //u_visible_light_indices[offset + i] = wg_visible_light_indices[i];
-        }
-
-        if (wg_visible_light_count != 1024) {
-            // Mark the end of the visible lights for this tile
-            u_visible_light_indices[offset + wg_visible_light_count] = -1;
-        }*/
     }
 
     workgroupBarrier();
 
     // now update the light index list on all threads.
-    var indices_offset = u32(index) * max_tile_visible_lights;
+    var indices_offset = u32(index) * MAX_TILE_VISIBLE_LIGHTS;
     //var pass_count = (wg_visible_light_count + thread_count - 1) / thread_count;
     for (var i = 0u; i < pass_count; i++) {
         // find the light index to check on this thread, make sure we're not trying to test
@@ -224,6 +209,20 @@ fn cs_main(
     }
 }
 
+/// 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 (offset < MAX_TILE_VISIBLE_LIGHTS) {
+        atomicAdd(&wg_visible_light_count, 1u);
+        wg_visible_light_indices[offset] = light_index;
+        return true;
+    }
+
+    return false;
+}
+
 fn sphere_inside_frustrum(frustum: array<vec4<f32>, 6>, sphere_origin: vec3<f32>, radius: f32) -> bool {
     // to be able to index this array with a non-const value,
     // it must be defined as a var
@@ -244,5 +243,7 @@ fn sphere_inside_frustrum(frustum: array<vec4<f32>, 6>, sphere_origin: vec3<f32>
 /// Source: Real-time collision detection, Christer Ericson (2005)
 ///   (https://www.3dgep.com/forward-plus/#light-culling-compute-shader)
 fn sphere_inside_plane(sphere_origin: vec3<f32>, radius: f32, plane: vec4<f32>) -> bool {
-    return dot(plane.xyz, sphere_origin) - plane.w < -radius;
+    //return dot(plane.xyz, sphere_origin) - plane.w < -radius;
+
+    return dot(vec4<f32>(sphere_origin, 0.0), plane) + radius > 0.0;
 }

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();