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.5,
+            look_speed: 0.3,
-            mouse_sensitivity: 0.9,
+            mouse_sensitivity: 1.0,
             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, 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::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::ResourceManager;
 
 mod free_fly_camera;
@@ -93,16 +93,15 @@ 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);
         
-        cube_gltf.wait_recurse_dependencies_load();
+        /* let cube_mesh = &cube_gltf.data_ref()
-        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()
@@ -152,33 +151,29 @@ async fn main() {
             ));
         }
 
-        {
+        /* {
-            //let mut light_tran = Transform::from_xyz(-3.5, 0.2, -4.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);
+            light_tran.scale = Vec3::new(0.5, 0.5, 0.5);
             world.spawn((
-                PointLight {
+                SpotLight {
-                    color: Vec3::new(0.0, 0.0, 1.0),
+                    color: Vec3::new(1.0, 0.2, 0.2),
-
+                    cutoff: math::Angle::Degrees(12.5),
-                    intensity: 3.3,
+                    outer_cutoff: math::Angle::Degrees(17.5),
 
                     constant: 1.0,
-                    linear: 0.09,
+                    linear: 0.007,
-                    quadratic: 0.032,
+                    quadratic: 0.0002,
 
-                    ambient: 0.2,
+                    ambient: 0.0,
-                    diffuse: 1.0,
+                    diffuse: 7.0,
-                    specular: 1.3,
+                    specular: 1.0,
                 },
-                Transform::new(
+                Transform::from(light_tran),
-                    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-engine/target/debug/testbed",
+        "executable": "/media/data_drive/Development/Rust/lyra-test/engine/target/debug/testbed",
         "inject": false,
         "numQueuedFrames": 1,
         "options": {
@@ -23,6 +23,6 @@
             "verifyBufferAccess": false
         },
         "queuedFrameCap": 5,
-        "workingDir": "/media/data_drive/Development/Rust/lyra-engine/examples/testbed"
+        "workingDir": "/media/data_drive/Development/Rust/lyra-test/engine/examples/testbed"
     }
 }

lyra-game/src/game.rs

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

lyra-game/src/render/light_cull_compute.rs

@@ -1,36 +1,47 @@
-use std::{borrow::Cow, mem, rc::Rc};
+use std::{borrow::Cow, mem, num::NonZeroU32, ptr::NonNull, rc::Rc};
 
 use glam::UVec2;
 use wgpu::{util::DeviceExt, ComputePipeline};
 use winit::dpi::PhysicalSize;
 
-use super::{light::LightUniformBuffers, render_buffer::{BindGroupPair, BufferWrapper}, texture::RenderTexture};
+use super::{light::LightUniformBuffers, render_buffer::BufferWrapper, renderer::RenderPass, texture::RenderTexture};
 
-pub(crate) struct LightIndicesGridBuffer {
+struct LightIndicesGridBuffer {
     indices_buffer: wgpu::Buffer,
     grid_texture: wgpu::Texture,
     grid_texture_view: wgpu::TextureView,
-    pub bg_pair: BindGroupPair,
+    bind_group_layout: wgpu::BindGroupLayout,
+    bind_group: wgpu::BindGroup,
 }
 
 pub(crate) struct LightCullCompute {
     device: Rc<wgpu::Device>,
     queue: Rc<wgpu::Queue>,
     pipeline: ComputePipeline,
-    pub light_indices_grid: LightIndicesGridBuffer,
+    lights: NonNull<LightUniformBuffers>,
+    camera: NonNull<BufferWrapper>,
+    light_indices_grid: LightIndicesGridBuffer,
     screen_size_buffer: BufferWrapper,
-    workgroup_size: glam::UVec2,
+    depth_tex: NonNull<RenderTexture>,
 }
 
 impl LightCullCompute {
-    fn create_grid(device: &wgpu::Device, screen_size: PhysicalSize<u32>, workgroup_size: glam::UVec2) -> LightIndicesGridBuffer {
+    fn create_grid(device: &wgpu::Device, screen_size: PhysicalSize<u32>) -> LightIndicesGridBuffer {
-        let mut contents = Vec::<u8>::new();
+        let limits = device.limits();
-        let contents_len = workgroup_size.x * workgroup_size.y * mem::size_of::<u8>() as u32;
+        //let max_buffer_sizes = (limits.max_uniform_buffer_binding_size as u64) / 2;
-        contents.resize(contents_len as _, 0);
+
+        /* let light_indices_buffer = device.create_buffer(
+            &wgpu::BufferDescriptor {
+                label: Some("B_LightIndices"),
+                usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
+                size: (mem::size_of::<u32>() * 16 * 16) as u64,
+                mapped_at_creation: false,
+            }
+        ); */
 
         let light_indices_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
             label: Some("B_LightIndices"),
-            contents: &contents,
+            contents: &[0; mem::size_of::<u32>() * 16 * 16],
             usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
         });
 
@@ -117,11 +128,12 @@ impl LightCullCompute {
             indices_buffer: light_indices_buffer,
             grid_texture,
             grid_texture_view,
-            bg_pair: BindGroupPair::new(light_indices_bg, light_indices_bg_layout),
+            bind_group_layout: light_indices_bg_layout,
+            bind_group: light_indices_bg,
         }
     }
 
-    pub fn new(device: Rc<wgpu::Device>, queue: Rc<wgpu::Queue>, screen_size: PhysicalSize<u32>, lights_buffers: &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: &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")
@@ -136,18 +148,18 @@ impl LightCullCompute {
             source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(shader_src)),
         });
 
-        let workgroup_size = glam::UVec2::new((screen_size.width as f32 / 16.0).ceil() as u32, (screen_size.height as f32 / 16.0).ceil() as u32);
+        let light_grid = Self::create_grid(&device, screen_size);
-        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(),
+                camera_buffers.bindgroup_layout().unwrap(),
-                &lights_buffers.bind_group_pair.layout,
+                &lights.bindgroup_layout,
-                &light_grid.bg_pair.layout,
+                &light_grid.bind_group_layout,
                 screen_size_buffer.bindgroup_layout().unwrap(),
             ],
             push_constant_ranges: &[],
@@ -164,19 +176,22 @@ impl LightCullCompute {
             device,
             queue,
             pipeline,
+            lights: NonNull::from(lights),
+            camera: NonNull::from(camera_buffers),
             light_indices_grid: light_grid,
             screen_size_buffer,
-            workgroup_size,
+            depth_tex: NonNull::from(depth_texture),
         }
     }
 
-    pub fn update_screen_size(&mut self, size: PhysicalSize<u32>) {
+    pub fn update_screen_size(&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, camera_buffers: &BufferWrapper, lights_buffers: &LightUniformBuffers, depth_texture: &RenderTexture) {
+    pub fn compute(&mut self) {
+        //self.queue.write_buffer(&self.light_indices_grid.indices_buffer, 0, &[0; mem::size_of::<u32>() * 16 * 16]);
+
         let mut encoder = self.device.create_command_encoder(&wgpu::CommandEncoderDescriptor {
             label: Some("LightCullCompute"),
         });
@@ -188,13 +203,19 @@ impl LightCullCompute {
 
             pass.set_pipeline(&self.pipeline);
             
-            pass.set_bind_group(0, depth_texture.bind_group(), &[]);
+            let depth = unsafe { self.depth_tex.as_ref() };
-            pass.set_bind_group(1, &camera_buffers.bindgroup(), &[]);
+            pass.set_bind_group(0, depth.bind_group(), &[]);
-            pass.set_bind_group(2, &lights_buffers.bind_group_pair.bindgroup, &[]);
+
-            pass.set_bind_group(3, &self.light_indices_grid.bg_pair.bindgroup, &[]);
+            let cam = unsafe { self.camera.as_ref() };
+            pass.set_bind_group(1, cam.bindgroup(), &[]);
+
+            let lights = unsafe { self.lights.as_ref() };
+            pass.set_bind_group(2, &lights.bindgroup, &[]);
+
+            pass.set_bind_group(3, &self.light_indices_grid.bind_group, &[]);
             pass.set_bind_group(4, self.screen_size_buffer.bindgroup(), &[]);
             
-            pass.dispatch_workgroups(self.workgroup_size.x, self.workgroup_size.y, 1);
+            pass.dispatch_workgroups(16, 16, 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::rc::Rc;
+use std::sync::Arc;
 
 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: &Rc<wgpu::BindGroupLayout>, i: &Option<ResHandle<Texture>>) -> Option<RenderTexture> {
+fn texture_to_render(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: &Arc<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: &Rc<
 }
 
 impl MaterialSpecular {
-    pub fn from_resource(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Rc<wgpu::BindGroupLayout>, value: &lyra_resource::gltf::Specular) -> Self {
+    pub fn from_resource(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Arc<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: Rc<wgpu::BindGroupLayout>, value: &lyra_resource::gltf::Material) -> Self {
+    pub fn from_resource(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Arc<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, rc::Rc};
+use std::{num::NonZeroU32, ops::Deref, sync::Arc};
 
 use wgpu::util::DeviceExt;
 
@@ -23,13 +23,13 @@ impl RenderBuffer {
 
 pub struct BindGroupPair {
     pub bindgroup: wgpu::BindGroup,
-    pub layout: Rc<wgpu::BindGroupLayout>,
+    pub layout: Arc<wgpu::BindGroupLayout>,
 }
 
 impl BindGroupPair {
-    pub fn create_bind_group(device: &wgpu::Device, layout: Rc<wgpu::BindGroupLayout>, entries: &[wgpu::BindGroupEntry<'_>]) -> Self {
+    pub fn new_from_layout(device: &wgpu::Device, layout: Arc<wgpu::BindGroupLayout>, entries: &[wgpu::BindGroupEntry<'_>]) -> Self {
         let bindgroup = device.create_bind_group(&wgpu::BindGroupDescriptor {
-            layout: &layout,
+            layout: layout.as_ref(),
             entries,
             label: None,
         });
@@ -43,7 +43,7 @@ impl BindGroupPair {
     pub fn new(bindgroup: wgpu::BindGroup, layout: wgpu::BindGroupLayout) -> Self {
         Self {
             bindgroup,
-            layout: Rc::new(layout),
+            layout: Arc::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)
+        self.bindgroup_pair.as_ref().map(|bg| bg.layout.deref())
     }
 
     /// 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: Rc::new(bg_layout),
+                    layout: bg_layout,
                 }
             }
         };

lyra-game/src/render/renderer.rs

@@ -59,7 +59,7 @@ struct MeshBufferStorage {
 
     //#[allow(dead_code)]
     //render_texture: Option<RenderTexture>,
-    material: Option<Rc<Material>>,
+    material: Option<Material>,
 
     // The index of the transform for this entity.
     // The tuple is structured like this: (transform index, index of transform inside the buffer)
@@ -88,7 +88,6 @@ 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>,
 
@@ -100,7 +99,7 @@ pub struct BasicRenderer {
     camera_buffer: BufferWrapper,
     //camera_bind_group: wgpu::BindGroup,
 
-    bgl_texture: Rc<BindGroupLayout>,
+    bgl_texture: Arc<BindGroupLayout>,
     default_texture: RenderTexture,
     depth_buffer_texture: RenderTexture,
 
@@ -173,7 +172,7 @@ impl BasicRenderer {
         };
         surface.configure(&device, &config);
 
-        let bgl_texture = Rc::new(RenderTexture::create_layout(&device));
+        let bgl_texture = Arc::new(RenderTexture::create_layout(&device));
 
         let shader_src = include_str!("shaders/base.wgsl");
         let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -224,7 +223,6 @@ impl BasicRenderer {
             render_pipelines: HashMap::new(),
             render_jobs: VecDeque::new(),
             mesh_buffers: HashMap::new(),
-            material_buffers: HashMap::new(),
             entity_meshes: HashMap::new(),
 
             render_limits,
@@ -249,10 +247,8 @@ 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.bind_group_pair.layout, &s.material_buffer.bindgroup_pair.as_ref().unwrap().layout,
+            &s.light_buffers.bindgroup_layout, &s.material_buffer.bindgroup_pair.as_ref().unwrap().layout,
-            &s.bgl_texture,
+            &s.bgl_texture])));
-            &s.light_cull_compute.light_indices_grid.bg_pair.layout,
-        ])));
         s.render_pipelines = pipelines;
 
         s
@@ -351,24 +347,18 @@ 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")
-        let material_ref = material.data_ref()
+            .data_ref()
             .unwrap();
-        
+        let material = Material::from_resource(&self.device, &self.queue, self.bgl_texture.clone(), &material);
-        let material = self.material_buffers.entry(material.uuid())
+        let uni = MaterialUniform::from(&material);
-            .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.clone()),
+            material: Some(material),
         }
     }
 
@@ -519,7 +509,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.camera_buffer, &self.light_buffers, &self.depth_buffer_texture);
+        //self.light_cull_compute.compute();
 
         let mut encoder = self.device.create_command_encoder(&wgpu::CommandEncoderDescriptor {
             label: Some("Basic Renderer's Encoder")
@@ -581,11 +571,9 @@ 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.bind_group_pair.bindgroup, &[]);
+                    render_pass.set_bind_group(3, &self.light_buffers.bindgroup, &[]);
                     render_pass.set_bind_group(4, &self.material_buffer.bindgroup_pair.as_ref().unwrap().bindgroup, &[]);
         
-                    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;
@@ -617,13 +605,7 @@ 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_directional = 0u;
-const LIGHT_TY_POINT = 1u;
+const light_ty_point = 1u;
-const LIGHT_TY_SPOT = 2u;
+const light_ty_spot = 2u;
 
 struct VertexInput {
     @location(0) position: vec3<f32>,
@@ -97,79 +97,41 @@ 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 block_size: i32 = 16;
-const MAX_TILE_VISIBLE_LIGHTS: u32 = 1024u;
+const max_tile_visible_lights: u32 = 1024u;
 
-const LIGHT_TY_DIRECTIONAL = 0u;
+const light_ty_directional = 0u;
-const LIGHT_TY_POINT = 1u;
+const light_ty_point = 1u;
-const LIGHT_TY_SPOT = 2u;
+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_depth_2d;
+var t_depthmap: texture_2d<f32>;
 @group(0) @binding(1)
 var s_depthmap: sampler;
 
@@ -62,8 +62,11 @@ 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>;
+var t_light_grid: texture_storage_2d<rg32uint, read_write>; // rg32uint = vec2<u32> or vec4<u32>(r, g, 0.0, 1.0) 
 
 @group(4) @binding(0)
 var<uniform> u_screen_size: vec2<u32>;
@@ -93,8 +96,9 @@ fn cs_main(
     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 tex_coord = vec2<f32>(global_invocation_id.xy);
-    var depth_float: f32 = textureLoad(t_depthmap, tex_coord, 0);
+    //var depth_float: f32 = textureSample(t_depthmap, s_depthmap, tex_coord).r;
+    var depth_float = 0.0;
     // bitcast the floating depth to u32 for atomic comparisons between threads
     var depth_uint: u32 = bitcast<u32>(depth_float);
 
@@ -142,7 +146,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
@@ -156,21 +160,18 @@ 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])) {
-                add_light(light_index);
+                var offset: u32 = wg_visible_light_count;
-                /*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;
-                }*/
+                }
             }
         }
     }
@@ -181,14 +182,28 @@ 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
+        var offset = u32(index) * max_tile_visible_lights; // index in the global light list
+        //t_light_grid[workgroup_id.x][workgroup_id.y] = vec2<f32>(offset, wg_visible_light_count);
         textureStore(t_light_grid, workgroup_id.xy, vec4<u32>(offset, wg_visible_light_count, 0u, 1u));
+        
+        // TODO: update transparent light grid
+
+        /*var offset = index * max_tile_visible_lights; // position in the global light buffer
+        // update the light
+        for (var i = 0u; i < wg_visible_light_count; i++) {
+            //u_visible_light_indices[offset + i] = wg_visible_light_indices[i];
+        }
+
+        if (wg_visible_light_count != 1024) {
+            // Mark the end of the visible lights for this tile
+            u_visible_light_indices[offset + wg_visible_light_count] = -1;
+        }*/
     }
 
     workgroupBarrier();
 
     // now update the light index list on all threads.
-    var indices_offset = u32(index) * MAX_TILE_VISIBLE_LIGHTS;
+    var 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
@@ -209,20 +224,6 @@ 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
@@ -243,7 +244,5 @@ 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::rc::Rc;
+use std::sync::Arc;
 
 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: Rc<wgpu::BindGroupLayout>, view: &wgpu::TextureView, sampler: &wgpu::Sampler) -> BindGroupPair {
+    fn create_bind_group_pair(device: &wgpu::Device, layout: Arc<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: Rc<wgpu::BindGroupLayout>, bytes: &[u8], label: &str) -> anyhow::Result<Self> {
+    pub fn from_bytes(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Arc<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: Rc<wgpu::BindGroupLayout>, img: &image::DynamicImage, label: Option<&str>) -> anyhow::Result<Self> {
+    pub fn from_image(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Arc<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: Rc<wgpu::BindGroupLayout>, texture_res: &ResHandle<Texture>, label: Option<&str>) -> anyhow::Result<Self> {
+    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> {
         let texture_ref = texture_res.data_ref().unwrap();
         let img = texture_ref.image.data_ref().unwrap();