From 22c08ba66e651c8573c70e8704eff627daaecabd Mon Sep 17 00:00:00 2001 From: SeanOMik Date: Thu, 14 Mar 2024 23:08:21 -0400 Subject: [PATCH 01/15] render: improve the render buffer wrapper, use it for the camera uniform --- lyra-game/src/render/camera.rs | 57 +++++++++++++++------- lyra-game/src/render/render_buffer.rs | 54 ++++++++++++++++++++- lyra-game/src/render/renderer.rs | 67 ++++++-------------------- lyra-game/src/render/shaders/base.wgsl | 3 +- 4 files changed, 111 insertions(+), 70 deletions(-) diff --git a/lyra-game/src/render/camera.rs b/lyra-game/src/render/camera.rs index f2f0729..07d73fc 100755 --- a/lyra-game/src/render/camera.rs +++ b/lyra-game/src/render/camera.rs @@ -1,3 +1,5 @@ +use std::{mem, num::NonZeroU64}; + use winit::dpi::PhysicalSize; use crate::{math::{Angle, OPENGL_TO_WGPU_MATRIX}, scene::CameraComponent}; @@ -38,16 +40,34 @@ impl Projection { #[repr(C)] #[derive(Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)] pub struct CameraUniform { + /// The view matrix of the camera + pub view_mat: glam::Mat4, + /// The view projection matrix pub view_proj: glam::Mat4, - // vec4 is used because of the uniforms 16 byte spacing requirement - pub view_pos: glam::Vec4, + /// The position of the camera + pub position: glam::Vec3, + _padding: u32, + } impl Default for CameraUniform { fn default() -> Self { Self { + view_mat: glam::Mat4::IDENTITY, view_proj: glam::Mat4::IDENTITY, - view_pos: Default::default() + position: Default::default(), + _padding: 0, + } + } +} + +impl CameraUniform { + pub fn new(view_mat: glam::Mat4, view_proj: glam::Mat4, position: glam::Vec3) -> Self { + Self { + view_mat, + view_proj, + position, + _padding: 0 } } } @@ -79,26 +99,29 @@ 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) -> (&glam::Mat4, glam::Mat4) { + 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) }, 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 +134,7 @@ 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 + (&self.view_proj, view) }, } } diff --git a/lyra-game/src/render/render_buffer.rs b/lyra-game/src/render/render_buffer.rs index 74d5e6d..2e92a52 100755 --- a/lyra-game/src/render/render_buffer.rs +++ b/lyra-game/src/render/render_buffer.rs @@ -1,4 +1,4 @@ -use std::{sync::Arc, num::NonZeroU32}; +use std::{num::NonZeroU32, ops::Deref, sync::Arc}; use wgpu::util::DeviceExt; @@ -68,11 +68,63 @@ impl BufferWrapper { } } + /// 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.deref()) + } + + /// 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(&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 bind_at_bind_group<'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); + } } +/// 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, diff --git a/lyra-game/src/render/renderer.rs b/lyra-game/src/render/renderer.rs index dda0383..8ded8da 100755 --- a/lyra-game/src/render/renderer.rs +++ b/lyra-game/src/render/renderer.rs @@ -88,8 +88,8 @@ pub struct BasicRenderer { render_limits: Limits, inuse_camera: RenderCamera, - camera_buffer: wgpu::Buffer, - camera_bind_group: wgpu::BindGroup, + camera_buffer: BufferWrapper, + //camera_bind_group: wgpu::BindGroup, bgl_texture: Arc, default_texture: RenderTexture, @@ -143,10 +143,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) { - true => wgpu::PresentMode::Immediate, - false => surface_caps.present_modes[0] - }; */ + let present_mode = surface_caps.present_modes[0]; debug!("present mode: {:?}", present_mode); @@ -174,41 +171,13 @@ impl BasicRenderer { }); let transform_buffers = TransformBuffers::new(&device); - - let camera_buffer = device.create_buffer_init( - &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 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); let depth_texture = RenderTexture::create_depth_texture(&device, &config, "Depth Buffer"); @@ -222,7 +191,6 @@ impl BasicRenderer { .buffer_usage(wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST) .visibility(wgpu::ShaderStages::FRAGMENT) .contents(&[MaterialUniform::default()]) - //.size(mem::size_of::()) .finish(&device); let mut s = Self { @@ -248,7 +216,6 @@ impl BasicRenderer { inuse_camera: RenderCamera::new(size), camera_buffer, - camera_bind_group, bgl_texture, default_texture, @@ -263,7 +230,8 @@ impl BasicRenderer { 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.camera_buffer.bindgroup_layout().unwrap(), &s.light_buffers.bindgroup_layout, &s.material_buffer.bindgroup_pair.as_ref().unwrap().layout, &s.bgl_texture]))); s.render_pipelines = pipelines; @@ -511,13 +479,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 (view_proj, view_mat) = self.inuse_camera.calc_view_projection(&camera); let pos = camera.transform.translation; - let uniform = CameraUniform { - view_proj: *view_proj, - 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])); + let uniform = CameraUniform::new(view_mat, *view_proj, pos); + self.camera_buffer.write_buffer(&self.queue, 0, &[uniform]); } else { warn!("Missing camera!"); } @@ -588,7 +553,7 @@ impl Renderer for BasicRenderer { let offset = TransformBuffers::index_offset(&self.render_limits, transform_indices) as u32; render_pass.set_bind_group(1, bindgroup, &[ offset, offset, ]); - render_pass.set_bind_group(2, &self.camera_bind_group, &[]); + self.camera_buffer.bind_at_bind_group(&mut render_pass, 2, &[]); 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, &[]); diff --git a/lyra-game/src/render/shaders/base.wgsl b/lyra-game/src/render/shaders/base.wgsl index dfbdc76..0930bf2 100755 --- a/lyra-game/src/render/shaders/base.wgsl +++ b/lyra-game/src/render/shaders/base.wgsl @@ -16,8 +16,9 @@ struct VertexOutput { } struct CameraUniform { + view_mat: mat4x4, view_proj: mat4x4, - view_pos: vec4, + view_pos: vec3, }; struct PointLight { -- 2.40.1 From cfd5cabfbb1eb68e7a8769c8b08d6e7b98cc084e Mon Sep 17 00:00:00 2001 From: SeanOMik Date: Sat, 16 Mar 2024 18:39:07 -0400 Subject: [PATCH 02/15] render: create light cull compute shader, bind buffers, etc. --- lyra-game/src/render/camera.rs | 2 - lyra-game/src/render/light/mod.rs | 189 ++++++++++--- lyra-game/src/render/light_cull_compute.rs | 223 ++++++++++++++++ lyra-game/src/render/mod.rs | 3 +- lyra-game/src/render/render_buffer.rs | 32 ++- lyra-game/src/render/renderer.rs | 30 ++- lyra-game/src/render/shaders/base.wgsl | 145 ++++------ .../src/render/shaders/light_cull.comp.wgsl | 248 ++++++++++++++++++ lyra-game/src/render/texture.rs | 60 +++++ 9 files changed, 782 insertions(+), 150 deletions(-) create mode 100644 lyra-game/src/render/light_cull_compute.rs create mode 100644 lyra-game/src/render/shaders/light_cull.comp.wgsl diff --git a/lyra-game/src/render/camera.rs b/lyra-game/src/render/camera.rs index 07d73fc..00d904d 100755 --- a/lyra-game/src/render/camera.rs +++ b/lyra-game/src/render/camera.rs @@ -1,5 +1,3 @@ -use std::{mem, num::NonZeroU64}; - use winit::dpi::PhysicalSize; use crate::{math::{Angle, OPENGL_TO_WGPU_MATRIX}, scene::CameraComponent}; diff --git a/lyra-game/src/render/light/mod.rs b/lyra-game/src/render/light/mod.rs index ee3c81b..9cdc09b 100644 --- a/lyra-game/src/render/light/mod.rs +++ b/lyra-game/src/render/light/mod.rs @@ -103,18 +103,21 @@ pub(crate) struct LightUniformBuffers { pub buffer: wgpu::Buffer, pub bindgroup_layout: wgpu::BindGroupLayout, pub bindgroup: wgpu::BindGroup, - pub lights_uniform: LightsUniform, - pub point_lights: LightBuffer, - pub spot_lights: LightBuffer, + pub light_indexes: HashMap, + pub current_light_idx: u32, } impl LightUniformBuffers { pub fn new(device: &wgpu::Device) -> Self { + let limits = device.limits(); + // TODO: check this limit somehow + let max_buffer_sizes = (limits.max_uniform_buffer_binding_size as u64) / 2; + let buffer = device.create_buffer( &wgpu::BufferDescriptor { label: Some("UBO_Lights"), - usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, - size: mem::size_of::() as u64, + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, + size: max_buffer_sizes, mapped_at_creation: false, } ); @@ -123,14 +126,16 @@ impl LightUniformBuffers { entries: &[ wgpu::BindGroupLayoutEntry { binding: 0, - visibility: wgpu::ShaderStages::VERTEX | wgpu::ShaderStages::FRAGMENT, + visibility: wgpu::ShaderStages::FRAGMENT | wgpu::ShaderStages::COMPUTE, ty: wgpu::BindingType::Buffer { - ty: wgpu::BufferBindingType::Uniform, + ty: wgpu::BufferBindingType::Storage { + read_only: true + }, has_dynamic_offset: false, min_binding_size: None, }, count: None, - } + }, ], label: Some("BGL_Lights"), }); @@ -147,63 +152,167 @@ impl LightUniformBuffers { size: None, // use the full buffer } ) - } + }, ], label: Some("BG_Lights"), }); - let point_lights = LightBuffer::new(MAX_LIGHT_COUNT); - let spot_lights = LightBuffer::new(MAX_LIGHT_COUNT); - Self { buffer, bindgroup_layout, bindgroup, - lights_uniform: LightsUniform::default(), - point_lights, - spot_lights, + light_indexes: Default::default(), + current_light_idx: 0, } } pub fn update_lights(&mut self, queue: &wgpu::Queue, world_tick: Tick, world: &World) { + let mut lights = LightsUniform::default(); + for (entity, point_light, transform, light_epoch, transform_epoch) in world.view_iter::<(Entities, &PointLight, &Transform, TickOf, TickOf)>() { - if !self.point_lights.has_light(entity) || light_epoch == world_tick || transform_epoch == world_tick { - let uniform = PointLightUniform::from_bundle(&point_light, &transform); - self.point_lights.update_or_add(&mut self.lights_uniform.point_lights, entity, uniform); - //debug!("Updated point light"); - } + // TODO: dont update light every frame + let idx = *self.light_indexes.entry(entity) + .or_insert_with(|| { + let t = self.current_light_idx; + self.current_light_idx += 1; + t + }) as usize; + + let uniform = LightUniform::from_point_light_bundle(&point_light, &transform); + lights.data[idx] = uniform; } - for (entity, spot_light, transform, light_epoch, transform_epoch) - in world.view_iter::<(Entities, &SpotLight, &Transform, TickOf, TickOf)>() { - - if !self.spot_lights.has_light(entity) || light_epoch == world_tick || transform_epoch == world_tick { - let uniform = SpotLightUniform::from_bundle(&spot_light, &transform); - self.spot_lights.update_or_add(&mut self.lights_uniform.spot_lights, entity, uniform); - //debug!("Updated spot light"); - } - } - - if let Some((dir_light, transform)) = - world.view_iter::<(&DirectionalLight, &Transform)>().next() { + if let Some((entity, dir_light, transform)) = + world.view_iter::<(Entities, &DirectionalLight, &Transform)>().next() { - let uniform = DirectionalLightUniform::from_bundle(&dir_light, &transform); - self.lights_uniform.directional_light = uniform; + let idx = *self.light_indexes.entry(entity) + .or_insert_with(|| { + let t = self.current_light_idx; + self.current_light_idx += 1; + t + }) as usize; + + let uniform = LightUniform::from_directional_bundle(&dir_light, &transform); + lights.data[idx] = uniform; } - self.lights_uniform.point_light_count = self.point_lights.buffer_count as u32; - self.lights_uniform.spot_light_count = self.spot_lights.buffer_count as u32; - queue.write_buffer(&self.buffer, 0, bytemuck::cast_slice(&[self.lights_uniform])); + lights.light_count = self.light_indexes.len() as u32; + + // update the light count in the struct + queue.write_buffer(&self.buffer, 0, bytemuck::cast_slice(&[lights])); } } +#[repr(C)] +#[derive(Default, Debug, Copy, Clone)] +pub(crate) enum LightType { + #[default] + Directional = 0, + Point = 1, + Spotlight = 2, +} + +#[repr(C)] +#[derive(Default, Debug, Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)] +pub(crate) struct LightUniform { + pub light_type: u32, // LightType + pub enabled: u32, // bool + pub _padding: [u32; 2], + pub position: glam::Vec3, + pub _padding2: u32, + pub direction: glam::Vec3, + pub _padding3: u32, + pub color: glam::Vec3, + // no padding is needed here since cutoff acts as the padding + // that would usually be needed for the vec3 + + pub range: f32, + pub intensity: f32, + + pub spot_cutoff: f32, + pub spot_outer_cutoff: f32, + + pub _padding4: u32, +} + +impl LightUniform { + pub fn from_point_light_bundle(light: &PointLight, transform: &Transform) -> Self { + Self { + light_type: LightType::Point as u32, + enabled: true as u32, // TODO + _padding: [0; 2], + position: transform.translation, + _padding2: 0, + direction: transform.forward(), + _padding3: 0, + color: light.color, + + range: 2.0, + intensity: 1.0, + + spot_cutoff: 0.0, + spot_outer_cutoff: 0.0, + + _padding4: 0, + } + } + + pub fn from_directional_bundle(light: &DirectionalLight, transform: &Transform) -> Self { + Self { + light_type: LightType::Directional as u32, + enabled: true as u32, // TODO: take from component + _padding: [0; 2], + position: transform.translation, + _padding2: 0, + direction: transform.forward(), + _padding3: 0, + color: light.color, + + range: 0.0, + intensity: 0.0, + + spot_cutoff: 0.0, + spot_outer_cutoff: 0.0, + + _padding4: 0, + } + } + + // Create the SpotLightUniform from an ECS bundle + /* pub fn from_bundle(light: &SpotLight, transform: &Transform) -> Self { + Self { + position: transform.translation, + _padding: 0, + direction: transform.forward(), + _padding2: 0, + color: light.color, + cutoff: light.cutoff.to_radians().cos(), + outer_cutoff: light.outer_cutoff.to_radians().cos(), + constant: light.constant, + linear: light.linear, + quadratic: light.quadratic, + ambient: light.ambient, + diffuse: light.diffuse, + specular: light.specular, + _padding3: 0, + } + } */ +} + #[repr(C)] #[derive(Default, Debug, Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)] pub struct LightsUniform { - point_lights: [PointLightUniform; MAX_LIGHT_COUNT], - point_light_count: u32, + light_count: u32, + _padding: [u32; 3], + data: [LightUniform; 10], // TODO: make this a dynamic length +} + +/* #[repr(C)] +#[derive(Default, Debug, Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)] +pub struct LightsUniform { + light_count: u32, _padding: [u32; 3], spot_lights: [SpotLightUniform; MAX_LIGHT_COUNT], spot_light_count: u32, @@ -352,4 +461,4 @@ impl SpotLightUniform { _padding3: 0, } } -} \ No newline at end of file +} */ \ No newline at end of file diff --git a/lyra-game/src/render/light_cull_compute.rs b/lyra-game/src/render/light_cull_compute.rs new file mode 100644 index 0000000..dca54c1 --- /dev/null +++ b/lyra-game/src/render/light_cull_compute.rs @@ -0,0 +1,223 @@ +use std::{borrow::Cow, mem, num::NonZeroU32, ptr::NonNull, rc::Rc}; + +use glam::UVec2; +use wgpu::{util::DeviceExt, ComputePipeline}; +use winit::dpi::PhysicalSize; + +use super::{light::LightUniformBuffers, render_buffer::BufferWrapper, renderer::RenderPass, texture::RenderTexture}; + +struct LightIndicesGridBuffer { + indices_buffer: wgpu::Buffer, + grid_texture: wgpu::Texture, + grid_texture_view: wgpu::TextureView, + bind_group_layout: wgpu::BindGroupLayout, + bind_group: wgpu::BindGroup, +} + +pub(crate) struct LightCullCompute { + device: Rc, + queue: Rc, + pipeline: ComputePipeline, + lights: NonNull, + camera: NonNull, + light_indices_grid: LightIndicesGridBuffer, + screen_size_buffer: BufferWrapper, + depth_tex: NonNull, +} + +impl LightCullCompute { + fn create_grid(device: &wgpu::Device, screen_size: PhysicalSize) -> 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::() * 16 * 16) as u64, + mapped_at_creation: false, + } + ); */ + + let light_indices_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor { + label: Some("B_LightIndices"), + contents: &[0; mem::size_of::() * 16 * 16], + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, + }); + + let light_indices_bg_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + entries: &[ + wgpu::BindGroupLayoutEntry { + binding: 0, + visibility: wgpu::ShaderStages::COMPUTE | wgpu::ShaderStages::FRAGMENT, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { + read_only: false + }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 1, + visibility: wgpu::ShaderStages::COMPUTE | wgpu::ShaderStages::FRAGMENT, + ty: wgpu::BindingType::StorageTexture { + access: wgpu::StorageTextureAccess::ReadWrite, + format: wgpu::TextureFormat::Rg32Uint, // vec2 + view_dimension: wgpu::TextureViewDimension::D2 + }, + count: None, + } + ], + label: Some("BGL_LightIndicesGrid"), + }); + + // TODO: shrink the texture to match the amount of grid cells that the shader actually uses + let size = wgpu::Extent3d { + width: screen_size.width, + height: screen_size.height, + depth_or_array_layers: 1, + }; + let grid_texture = device.create_texture( + &wgpu::TextureDescriptor { + label: Some("Tex_LightGrid"), + size, + mip_level_count: 1, + sample_count: 1, + dimension: wgpu::TextureDimension::D2, + format: wgpu::TextureFormat::Rg32Uint, // vec2 + 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 + dimension: Some(wgpu::TextureViewDimension::D2), + aspect: wgpu::TextureAspect::All, + base_mip_level: 0, + mip_level_count: None, + base_array_layer: 0, + array_layer_count: None, + }); + + let light_indices_bg = device.create_bind_group(&wgpu::BindGroupDescriptor { + layout: &light_indices_bg_layout, + entries: &[ + wgpu::BindGroupEntry { + binding: 0, + resource: wgpu::BindingResource::Buffer( + wgpu::BufferBinding { + buffer: &light_indices_buffer, + offset: 0, + size: None, // the entire light buffer is needed + } + ) + }, + wgpu::BindGroupEntry { + binding: 1, + resource: wgpu::BindingResource::TextureView(&grid_texture_view) + } + ], + label: Some("BG_LightIndicesGrid"), + }); + + LightIndicesGridBuffer { + indices_buffer: light_indices_buffer, + grid_texture, + grid_texture_view, + bind_group_layout: light_indices_bg_layout, + bind_group: light_indices_bg, + } + } + + pub fn new(device: Rc, queue: Rc, screen_size: PhysicalSize, lights: &LightUniformBuffers, camera_buffers: &BufferWrapper, depth_texture: &mut RenderTexture) -> Self { + let screen_size_buffer = BufferWrapper::builder() + .buffer_usage(wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST) + .label_prefix("ScreenSize") + .visibility(wgpu::ShaderStages::COMPUTE) + .buffer_dynamic_offset(false) + .contents(&[UVec2::new(screen_size.width, screen_size.height)]) + .finish(&device); + + let shader_src = include_str!("shaders/light_cull.comp.wgsl"); + let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor { + label: Some("LightCullCompute"), + source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(shader_src)), + }); + + let light_grid = Self::create_grid(&device, screen_size); + + let depth_tex_pair = depth_texture.create_bind_group(&device); + + let layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { + label: Some("PipeLay_LightCull"), + bind_group_layouts: &[ + //&depth_texture.bindgroup_pair.as_ref().unwrap().layout, + &depth_tex_pair.layout, + camera_buffers.bindgroup_layout().unwrap(), + &lights.bindgroup_layout, + &light_grid.bind_group_layout, + screen_size_buffer.bindgroup_layout().unwrap(), + ], + push_constant_ranges: &[], + }); + + let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: Some("Pipe_LightCull"), + layout: Some(&layout), + module: &shader, + entry_point: "cs_main", + }); + + Self { + device, + queue, + pipeline, + lights: NonNull::from(lights), + camera: NonNull::from(camera_buffers), + light_indices_grid: light_grid, + screen_size_buffer, + depth_tex: NonNull::from(depth_texture), + } + } + + pub fn update_screen_size(&self, size: PhysicalSize) { + self.screen_size_buffer.write_buffer(&self.queue, 0, + &[UVec2::new(size.width, size.height)]); + } + + pub fn compute(&mut self) { + //self.queue.write_buffer(&self.light_indices_grid.indices_buffer, 0, &[0; mem::size_of::() * 16 * 16]); + + let mut encoder = self.device.create_command_encoder(&wgpu::CommandEncoderDescriptor { + label: Some("LightCullCompute"), + }); + + { + let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: Some("Pass_LightCull"), + }); + + pass.set_pipeline(&self.pipeline); + + let depth = unsafe { self.depth_tex.as_ref() }; + pass.set_bind_group(0, depth.bind_group(), &[]); + + let cam = unsafe { self.camera.as_ref() }; + pass.set_bind_group(1, cam.bindgroup(), &[]); + + let lights = unsafe { self.lights.as_ref() }; + pass.set_bind_group(2, &lights.bindgroup, &[]); + + pass.set_bind_group(3, &self.light_indices_grid.bind_group, &[]); + pass.set_bind_group(4, self.screen_size_buffer.bindgroup(), &[]); + + pass.dispatch_workgroups(16, 16, 1); + } + self.queue.submit(std::iter::once(encoder.finish())); + self.device.poll(wgpu::Maintain::Wait); + } +} \ No newline at end of file diff --git a/lyra-game/src/render/mod.rs b/lyra-game/src/render/mod.rs index 0830bba..5475230 100755 --- a/lyra-game/src/render/mod.rs +++ b/lyra-game/src/render/mod.rs @@ -11,4 +11,5 @@ pub mod material; pub mod camera; pub mod window; pub mod transform_buffer_storage; -pub mod light; \ No newline at end of file +pub mod light; +pub mod light_cull_compute; \ No newline at end of file diff --git a/lyra-game/src/render/render_buffer.rs b/lyra-game/src/render/render_buffer.rs index 2e92a52..df174e0 100755 --- a/lyra-game/src/render/render_buffer.rs +++ b/lyra-game/src/render/render_buffer.rs @@ -39,12 +39,19 @@ impl BindGroupPair { layout, } } + + pub fn new(bindgroup: wgpu::BindGroup, layout: wgpu::BindGroupLayout) -> Self { + Self { + bindgroup, + layout: Arc::new(layout), + } + } } pub struct BufferWrapper { pub bindgroup_pair: Option, pub inner_buf: wgpu::Buffer, - pub len: usize, + pub len: Option, } impl BufferWrapper { @@ -54,7 +61,7 @@ impl BufferWrapper { Self { bindgroup_pair: bind_group, inner_buf: buffer, - len: 0, + len: Some(0), } } @@ -64,7 +71,15 @@ impl BufferWrapper { Self { bindgroup_pair: bind_group, inner_buf: buffer, - len: 0, + len: Some(0), + } + } + + pub fn from_parts(bind_group: wgpu::BindGroup, bind_group_layout: wgpu::BindGroupLayout, buffer: wgpu::Buffer) -> Self { + Self { + bindgroup_pair: Some(BindGroupPair::new(bind_group, bind_group_layout)), + inner_buf: buffer, + len: None, } } @@ -101,7 +116,7 @@ impl BufferWrapper { /// match the layout of this bind group. /// /// See [`wgpu::RenderPass::set_bind_group`](https://docs.rs/wgpu/latest/wgpu/struct.RenderPass.html#method.set_bind_group). - pub fn bind_at_bind_group<'a, 'b>( + pub fn render_pass_bind_at<'a, 'b>( &'a self, pass: &'b mut wgpu::RenderPass<'a>, index: u32, @@ -112,6 +127,13 @@ impl BufferWrapper { ); pass.set_bind_group(index, &pair.bindgroup, offsets); } + + /// Returns the bindgroup of this buffer, panics if the buffer does not have a bindgroup. + pub fn bindgroup(&self) -> &wgpu::BindGroup { + &self.bindgroup_pair.as_ref().expect( + "BufferWrapper is missing bindgroup pair! Cannot set bind group on RenderPass!", + ).bindgroup + } } /// Struct used for building a BufferWrapper @@ -274,7 +296,7 @@ impl BufferWrapperBuilder { BufferWrapper { bindgroup_pair: Some(bg_pair), inner_buf: buffer, - len: self.count.unwrap_or_default() as usize, + len: Some(self.count.unwrap_or_default() as usize), } } } diff --git a/lyra-game/src/render/renderer.rs b/lyra-game/src/render/renderer.rs index 8ded8da..47eda68 100755 --- a/lyra-game/src/render/renderer.rs +++ b/lyra-game/src/render/renderer.rs @@ -1,4 +1,5 @@ use std::collections::{HashMap, VecDeque, HashSet}; +use std::rc::Rc; use std::sync::Arc; use std::borrow::Cow; @@ -24,6 +25,7 @@ use crate::scene::CameraComponent; use super::camera::{RenderCamera, CameraUniform}; use super::desc_buf_lay::DescVertexBufferLayout; use super::light::LightUniformBuffers; +use super::light_cull_compute::LightCullCompute; use super::material::Material; use super::render_buffer::BufferWrapper; use super::texture::RenderTexture; @@ -45,6 +47,12 @@ pub trait Renderer { fn add_render_pipeline(&mut self, shader_id: u64, pipeline: Arc); } +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); +} + struct MeshBufferStorage { buffer_vertex: BufferStorage, buffer_indices: Option<(wgpu::IndexFormat, BufferStorage)>, @@ -68,8 +76,8 @@ pub struct CachedTransform { pub struct BasicRenderer { pub surface: wgpu::Surface, - pub device: wgpu::Device, - pub queue: wgpu::Queue, + pub device: Rc, // device does not need to be mutable, no need for refcell + pub queue: Rc, pub config: wgpu::SurfaceConfiguration, pub size: winit::dpi::PhysicalSize, pub window: Arc, @@ -98,6 +106,8 @@ pub struct BasicRenderer { material_buffer: BufferWrapper, light_buffers: LightUniformBuffers, + + light_cull_compute: LightCullCompute, } impl BasicRenderer { @@ -123,7 +133,7 @@ impl BasicRenderer { let (device, queue) = adapter.request_device( &wgpu::DeviceDescriptor { - features: wgpu::Features::empty(), + features: wgpu::Features::TEXTURE_ADAPTER_SPECIFIC_FORMAT_FEATURES, // WebGL does not support all wgpu features. // Not sure if the engine will ever completely support WASM, // but its here just in case @@ -175,11 +185,11 @@ impl BasicRenderer { .buffer_usage(wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST) .contents(&[CameraUniform::default()]) .label_prefix("Camera") - .visibility(wgpu::ShaderStages::VERTEX | wgpu::ShaderStages::FRAGMENT) + .visibility(wgpu::ShaderStages::VERTEX | wgpu::ShaderStages::FRAGMENT | wgpu::ShaderStages::COMPUTE) .buffer_dynamic_offset(false) .finish(&device); - let depth_texture = RenderTexture::create_depth_texture(&device, &config, "Depth Buffer"); + let mut depth_texture = RenderTexture::create_depth_texture(&device, &config, "Tex_Depth"); // load the default texture let bytes = include_bytes!("default_texture.png"); @@ -193,6 +203,10 @@ impl BasicRenderer { .contents(&[MaterialUniform::default()]) .finish(&device); + let device = Rc::new(device); + let queue = Rc::new(queue); + let light_cull_compute = LightCullCompute::new(device.clone(), queue.clone(), size, &light_uniform_buffers, &camera_buffer, &mut depth_texture); + let mut s = Self { window, surface, @@ -224,6 +238,7 @@ impl BasicRenderer { light_buffers: light_uniform_buffers, material_buffer: mat_buffer, + light_cull_compute, }; // create the default pipelines @@ -494,6 +509,8 @@ impl Renderer for BasicRenderer { let output = self.surface.get_current_texture()?; let view = output.texture.create_view(&wgpu::TextureViewDescriptor::default()); + //self.light_cull_compute.compute(); + let mut encoder = self.device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: Some("Basic Renderer's Encoder") }); @@ -553,7 +570,7 @@ impl Renderer for BasicRenderer { let offset = TransformBuffers::index_offset(&self.render_limits, transform_indices) as u32; render_pass.set_bind_group(1, bindgroup, &[ offset, offset, ]); - self.camera_buffer.bind_at_bind_group(&mut render_pass, 2, &[]); + render_pass.set_bind_group(2, &self.camera_buffer.bindgroup(), &[]); render_pass.set_bind_group(3, &self.light_buffers.bindgroup, &[]); render_pass.set_bind_group(4, &self.material_buffer.bindgroup_pair.as_ref().unwrap().bindgroup, &[]); @@ -590,6 +607,7 @@ impl Renderer for BasicRenderer { self.surface.configure(&self.device, &self.config); self.depth_buffer_texture = RenderTexture::create_depth_texture(&self.device, &self.config, "Depth Buffer Texture"); self.inuse_camera.update_aspect_ratio(self.size); + self.light_cull_compute.update_screen_size(new_size); } } diff --git a/lyra-game/src/render/shaders/base.wgsl b/lyra-game/src/render/shaders/base.wgsl index 0930bf2..c3c8bc1 100755 --- a/lyra-game/src/render/shaders/base.wgsl +++ b/lyra-game/src/render/shaders/base.wgsl @@ -2,6 +2,10 @@ const max_light_count: u32 = 16u; +const light_ty_directional = 0u; +const light_ty_point = 1u; +const light_ty_spot = 2u; + struct VertexInput { @location(0) position: vec3, @location(1) tex_coords: vec2, @@ -21,53 +25,25 @@ struct CameraUniform { view_pos: vec3, }; -struct PointLight { - position: vec4, - color: vec4, - - intensity: f32, - constant: f32, - linear: f32, - quadratic: f32, - - ambient: f32, - diffuse: f32, - specular: f32, -}; +struct Light { + light_ty: u32, + enabled: u32, -struct DirectionalLight { - direction: vec3, - color: vec3, - - ambient: f32, - diffuse: f32, - specular: f32, -}; - -struct SpotLight { position: vec3, direction: vec3, color: vec3, - cutoff: f32, - outer_cutoff: f32, + range: f32, + intensity: f32, - constant: f32, - linear: f32, - quadratic: f32, - - ambient: f32, - diffuse: f32, - specular: f32, + spot_cutoff: f32, + spot_outer_cutoff: f32, }; struct Lights { - point_lights: array, - point_light_count: u32, - spot_lights: array, - spot_light_count: u32, - directional_light: DirectionalLight, -} + light_count: u32, + data: array, +}; @group(1) @binding(0) var u_model_transform: mat4x4; @@ -78,7 +54,7 @@ var u_model_normal_matrix: mat4x4; var u_camera: CameraUniform; @group(3) @binding(0) -var u_lights: Lights; +var u_lights: Lights; @vertex fn vs_main( @@ -129,7 +105,21 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4 { // this needs to be 0.0 for the math //u_lights.directional_light.direction.w = 0.0; - var light_res = blinn_phong_dir_light(in.world_position, in.world_normal, u_lights.directional_light, u_material, specular_color); + var light_res = vec3(0.0); + + for (var i = 0u; i < u_lights.light_count; i++) { + var light = u_lights.data[i]; + + if (light.light_ty == light_ty_directional) { + light_res += blinn_phong_dir_light(in.world_position, in.world_normal, light, u_material, specular_color); + } else if (light.light_ty == light_ty_point) { + light_res += blinn_phong_point_light(in.world_position, in.world_normal, light, u_material, specular_color); + } else if (light.light_ty == light_ty_spot) { + light_res += blinn_phong_spot_light(in.world_position, in.world_normal, light, u_material, specular_color); + } + } + + /*var light_res = blinn_phong_dir_light(in.world_position, in.world_normal, u_lights.directional_light, u_material, specular_color); for (var i = 0u; i < u_lights.point_light_count; i++) { light_res += blinn_phong_point_light(in.world_position, in.world_normal, u_lights.point_lights[i], u_material, specular_color); @@ -137,14 +127,14 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4 { 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(light_object_res, object_color.a); } -fn blinn_phong_dir_light(world_pos: vec3, world_norm: vec3, dir_light: DirectionalLight, material: Material, specular_factor: vec3) -> vec3 { +fn blinn_phong_dir_light(world_pos: vec3, world_norm: vec3, dir_light: Light, material: Material, specular_factor: vec3) -> vec3 { let light_color = dir_light.color.xyz; let camera_view_pos = u_camera.view_pos.xyz; @@ -166,14 +156,14 @@ fn blinn_phong_dir_light(world_pos: vec3, world_norm: vec3, dir_light: var specular_color = specular_strength * (light_color * specular_factor); //// end of specular //// - ambient_color *= dir_light.ambient; + /*ambient_color *= dir_light.ambient; diffuse_color *= dir_light.diffuse; - specular_color *= dir_light.specular; + specular_color *= dir_light.specular;*/ return ambient_color + diffuse_color + specular_color; } -fn blinn_phong_point_light(world_pos: vec3, world_norm: vec3, point_light: PointLight, material: Material, specular_factor: vec3) -> vec3 { +fn blinn_phong_point_light(world_pos: vec3, world_norm: vec3, point_light: Light, material: Material, specular_factor: vec3) -> vec3 { let light_color = point_light.color.xyz; let light_pos = point_light.position.xyz; let camera_view_pos = u_camera.view_pos.xyz; @@ -196,8 +186,14 @@ fn blinn_phong_point_light(world_pos: vec3, world_norm: vec3, point_li var specular_color = specular_strength * (light_color * specular_factor); //// end of specular //// - //// point light attenuation //// + // TODO: Point light range let distance = length(light_pos - world_pos); + // TODO: make smoothness in this a configurable value + // 0.75 is the smoothness or falloff + let attenuation = 1.0 - smoothstep(point_light.range * 0.75, point_light.range, distance); + + //// point light attenuation //// + /*let distance = length(light_pos - world_pos); let attenuation = 1.0 / (point_light.constant + point_light.linear * distance + point_light.quadratic * (distance * distance)); @@ -205,58 +201,15 @@ fn blinn_phong_point_light(world_pos: vec3, world_norm: vec3, point_li ambient_color *= point_light.ambient * attenuation; diffuse_color *= point_light.diffuse * attenuation; - specular_color *= point_light.specular * attenuation; + specular_color *= point_light.specular * attenuation;*/ + + ambient_color *= attenuation; + diffuse_color *= attenuation; + specular_color *= attenuation; return (ambient_color + diffuse_color + specular_color) * point_light.intensity; } -fn blinn_phong_spot_light(world_pos: vec3, world_norm: vec3, spot_light: SpotLight, material: Material, specular_factor: vec3) -> vec3 { - let light_color = spot_light.color;//.xyz; - let light_pos = spot_light.position.xyz; - let camera_view_pos = u_camera.view_pos.xyz; - - let light_dir = normalize(spot_light.position - world_pos); - - - //if (theta > spot_light.cutoff) { - var ambient_color = light_color * material.ambient.xyz * material.diffuse.xyz; - - //// diffuse //// - //let light_dir = normalize(light_pos - world_pos); - - let diffuse_strength = max(dot(world_norm, light_dir), 0.0); - var diffuse_color = light_color * (diffuse_strength * material.diffuse.xyz); - //// end of diffuse //// - - //// specular //// - let view_dir = normalize(camera_view_pos - world_pos); - let half_dir = normalize(view_dir + light_dir); - - let specular_strength = pow(max(dot(world_norm, half_dir), 0.0), material.shininess); - var specular_color = specular_strength * (light_color * specular_factor); - //// end of specular //// - - //// spot light soft edges //// - let theta = dot(light_dir, normalize(-spot_light.direction)); - let epsilon = spot_light.cutoff - spot_light.outer_cutoff; - let intensity = clamp((theta - spot_light.outer_cutoff) / epsilon, 0.0, 1.0); - //diffuse_color *= intensity; - //specular_color *= intensity; - //// end of spot light soft edges //// - - //// spot light attenuation //// - let distance = length(light_pos - world_pos); - let attenuation = 1.0 / (spot_light.constant + spot_light.linear * distance + - spot_light.quadratic * (distance * distance)); - - ambient_color *= attenuation * intensity * spot_light.ambient; - diffuse_color *= attenuation * intensity * spot_light.diffuse; - specular_color *= attenuation * intensity * spot_light.specular; - //// end of spot light attenuation //// - - - return /*ambient_color +*/ diffuse_color + specular_color; - /*} else { - return vec3(0.0); - }*/ +fn blinn_phong_spot_light(world_pos: vec3, world_norm: vec3, spot_light: Light, material: Material, specular_factor: vec3) -> vec3 { + return vec3(0.0); // TODO } \ No newline at end of file diff --git a/lyra-game/src/render/shaders/light_cull.comp.wgsl b/lyra-game/src/render/shaders/light_cull.comp.wgsl new file mode 100644 index 0000000..54df3c5 --- /dev/null +++ b/lyra-game/src/render/shaders/light_cull.comp.wgsl @@ -0,0 +1,248 @@ +const block_size: i32 = 16; +const max_tile_visible_lights: u32 = 1024u; + +const light_ty_directional = 0u; +const light_ty_point = 1u; +const light_ty_spot = 2u; + +// Possible computer shader inputs: +// +// local_invocation_id +// workgroup_id +// global_invocation_id +// num_workgroups +// local_invocation_index + +struct CameraUniform { + view_mat: mat4x4, + view_proj: mat4x4, + view_pos: vec3, +}; + +struct Light { + light_ty: u32, + enabled: u32, + + position: vec3, + direction: vec3, + color: vec3, + + range: f32, + intensity: f32, + + spot_cutoff: f32, + spot_outer_cutoff: f32, +}; + +struct Lights { + light_count: u32, + data: array, +}; + +var wg_min_depth: atomic; +var wg_max_depth: atomic; +var wg_frustum_planes: array, 6>; + +// index list of visible light sources for this tile +var wg_visible_light_indices: array; +var wg_visible_light_count: atomic; + +//var view_projection: mat4x4; + +@group(0) @binding(0) +var t_depthmap: texture_2d; +@group(0) @binding(1) +var s_depthmap: sampler; + +@group(1) @binding(0) +var u_camera: CameraUniform; + +@group(2) @binding(0) +var u_lights: Lights; + +@group(3) @binding(0) +var u_light_indices: array; +/*@group(3) @binding(1) +var u_light_grid: array>>;*/ + +@group(3) @binding(1) +var t_light_grid: texture_storage_2d; // rg32uint = vec2 or vec4(r, g, 0.0, 1.0) + +@group(4) @binding(0) +var u_screen_size: vec2; + +@compute +@workgroup_size(16, 16, 1) +fn cs_main( + @builtin(local_invocation_id) local_invocation_id: vec3, + @builtin(workgroup_id) workgroup_id: vec3, + @builtin(global_invocation_id) global_invocation_id: vec3, + @builtin(num_workgroups) num_workgroups: vec3, + @builtin(local_invocation_index) local_invocation_index: u32, +) { + //var location = vec2(global_invocation_id.xy); + var item_id = vec2(local_invocation_id.xy); + var tile_id = vec2(workgroup_id.xy); + var tile_number = vec2(num_workgroups.xy); + var index = tile_id.y * tile_number.x + tile_id.x; + + // Initialize some shared global values for depth and light count + if (local_invocation_index == 0u) { + wg_min_depth = 0xFFFFFFFu; + wg_max_depth = 0u; + wg_visible_light_count = 0u; + } + + workgroupBarrier(); + + // step 1: calculate the minimum and maximum depth values for this tile (using the depth map) + var tex_coord = vec2(global_invocation_id.xy); + //var depth_float: f32 = textureSample(t_depthmap, s_depthmap, tex_coord).r; + var depth_float = 0.0; + // bitcast the floating depth to u32 for atomic comparisons between threads + var depth_uint: u32 = bitcast(depth_float); + + // step 2: find the minimum and max depth for this tile. + // atomically update the workgroup depth + atomicMin(&wg_min_depth, depth_uint); + atomicMax(&wg_max_depth, depth_uint); + + // convert them back into floats + var min_depth: f32 = bitcast(wg_min_depth); + var max_depth: f32 = bitcast(wg_max_depth); + + workgroupBarrier(); + + // Create the frustum planes that will be used for this time + if (local_invocation_index == 0u) { + var negative_step = (2.0 * vec2(tile_id)) / vec2(tile_number); + var positive_step = (2.0 * vec2(tile_id) + vec2(1.0, 1.0)) / vec2(tile_number); + + // z in the vec4 is the distance from the center of the tile + wg_frustum_planes[0] = vec4(1.0, 0.0, 0.0, 1.0 - negative_step.x); // left + wg_frustum_planes[1] = vec4(-1.0, 0.0, 0.0, -1.0 + positive_step.x); // right + wg_frustum_planes[2] = vec4(0.0, -1.0, 0.0, 1.0 - negative_step.y); // bottom + wg_frustum_planes[3] = vec4(0.0, -1.0, 0.0, -1.0 + positive_step.y); // top + wg_frustum_planes[4] = vec4(0.0, 0.0, -1.0, -min_depth); // near plane + wg_frustum_planes[5] = vec4(0.0, 0.0, 1.0, max_depth); // far plane + + // convert the side and top planes from clip to view space + for (var i = 0u; i < 4u; i++) { + wg_frustum_planes[i] *= u_camera.view_proj; + wg_frustum_planes[i] /= length(wg_frustum_planes[i].xyz); + } + + // convert near and far planes from clip to view space + wg_frustum_planes[4] *= u_camera.view_mat; + wg_frustum_planes[4] /= length(wg_frustum_planes[4].xyz); + wg_frustum_planes[5] *= u_camera.view_mat; + wg_frustum_planes[5] /= length(wg_frustum_planes[5].xyz); + } + + workgroupBarrier(); + + // Step 3: cull lights + + // Process the lights detecting which ones to cull for this tile. + // Processes 256 lights simultaniously, each on a thread in the workgroup. Requires multiple + // iterations for more lights. + var thread_count = u32(block_size * block_size); + var pass_count = (u_lights.light_count + thread_count - 1u) / thread_count; + for (var i = 0u; i < pass_count; i++) { + // find the light index to check on this thread, make sure we're not trying to test + // for more lights than we have. + var light_index = i * thread_count + local_invocation_index; + if (light_index > u_lights.light_count) { + break; + } + + var light = u_lights.data[light_index]; + var position = light.position; + var radius = light.range; + + if (light.light_ty != light_ty_directional + && sphere_inside_frustrum(wg_frustum_planes, position, radius)) { + // TODO: add the light to the transparent geometry list + + // TODO: spotlights + if (!sphere_inside_plane(position, radius, wg_frustum_planes[4])) { + var offset: u32 = wg_visible_light_count; + + if (offset < max_tile_visible_lights) { + atomicAdd(&wg_visible_light_count, 1u); + wg_visible_light_indices[offset] = light_index; + } + } + } + } + + workgroupBarrier(); + + // Update the global memory with the visible light buffer. + + // first update the light grid on the first thread + if (local_invocation_index == 0u) { + var offset = u32(index) * max_tile_visible_lights; // index in the global light list + //t_light_grid[workgroup_id.x][workgroup_id.y] = vec2(offset, wg_visible_light_count); + textureStore(t_light_grid, workgroup_id.xy, vec4(offset, wg_visible_light_count, 0u, 1u)); + + // TODO: update transparent light grid + + /*var offset = index * max_tile_visible_lights; // position in the global light buffer + // update the light + for (var i = 0u; i < wg_visible_light_count; i++) { + //u_visible_light_indices[offset + i] = wg_visible_light_indices[i]; + } + + if (wg_visible_light_count != 1024) { + // Mark the end of the visible lights for this tile + u_visible_light_indices[offset + wg_visible_light_count] = -1; + }*/ + } + + workgroupBarrier(); + + // now update the light index list on all threads. + var indices_offset = u32(index) * max_tile_visible_lights; + //var pass_count = (wg_visible_light_count + thread_count - 1) / thread_count; + for (var i = 0u; i < pass_count; i++) { + // find the light index to check on this thread, make sure we're not trying to test + // for more lights than we have. + //var light_index: u32 = i * thread_count + local_invocation_index; + /*if (light_index > u_lights.light_count) { + u_visible_light_indices + break; + }*/ + + var offset = indices_offset + i; + if (offset >= wg_visible_light_count) { + // stop if we're over the over the amount of lights we saw + break; + } + + u_light_indices[offset] = wg_visible_light_indices[i]; + } +} + +fn sphere_inside_frustrum(frustum: array, 6>, sphere_origin: vec3, radius: f32) -> bool { + // to be able to index this array with a non-const value, + // it must be defined as a var + var frustum_v = frustum; + + // only check the sides of the frustum + for (var i = 0u; i < 4u; i++) { + if (!sphere_inside_plane(sphere_origin, radius, frustum_v[i])) { + return false; + } + } + + return true; +} + +/// Check if the sphere is fully behind (i.e., inside the negative half-space of) a plane. +/// +/// Source: Real-time collision detection, Christer Ericson (2005) +/// (https://www.3dgep.com/forward-plus/#light-culling-compute-shader) +fn sphere_inside_plane(sphere_origin: vec3, radius: f32, plane: vec4) -> bool { + return dot(plane.xyz, sphere_origin) - plane.w < -radius; +} \ No newline at end of file diff --git a/lyra-game/src/render/texture.rs b/lyra-game/src/render/texture.rs index e17a0ab..154ad11 100755 --- a/lyra-game/src/render/texture.rs +++ b/lyra-game/src/render/texture.rs @@ -297,6 +297,66 @@ impl RenderTexture { } } + /// Creates a bind group for this texture and returns a borrow to the [`BindGroupPair`] + /// + /// This does not create a new bind group if the texture already has one. + /// The view dimension will be the same as the texture dimension. + pub fn create_bind_group(&mut self, device: &wgpu::Device) -> &BindGroupPair { + if self.bindgroup_pair.is_some() { + // could not use an if-let here due to the borrow checker thinking + // that there was multiple borrows to self.bindgroup_pair + return self.bindgroup_pair.as_ref().unwrap(); + } + + let view_dim = match self.inner_texture.dimension() { + wgpu::TextureDimension::D1 => wgpu::TextureViewDimension::D1, + wgpu::TextureDimension::D2 => wgpu::TextureViewDimension::D2, + wgpu::TextureDimension::D3 => wgpu::TextureViewDimension::D3, + }; + + let layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + entries: &[ + wgpu::BindGroupLayoutEntry { + binding: 0, + visibility: wgpu::ShaderStages::COMPUTE | wgpu::ShaderStages::FRAGMENT, + ty: wgpu::BindingType::Texture { + sample_type: wgpu::TextureSampleType::Depth, + view_dimension: view_dim, + multisampled: false + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 1, + visibility: wgpu::ShaderStages::COMPUTE | wgpu::ShaderStages::FRAGMENT, + ty: wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Comparison), + count: None, + } + ], + label: Some("BGL_Texture"), + }); + + let bg = device.create_bind_group(&wgpu::BindGroupDescriptor { + layout: &layout, + entries: &[ + wgpu::BindGroupEntry { + binding: 0, + resource: wgpu::BindingResource::TextureView(&self.view) + }, + wgpu::BindGroupEntry { + binding: 1, + resource: wgpu::BindingResource::Sampler(&self.sampler) + } + ], + label: Some("BG_Texture"), + }); + + let pair = BindGroupPair::new(bg, layout); + + self.bindgroup_pair = Some(pair); + self.bindgroup_pair.as_ref().unwrap() + } + /// Returns the bind group stored inside the bind group pair. /// /// Panics: -- 2.40.1 From 1818a0b48b33a14170cdc867f64c82d8df52c265 Mon Sep 17 00:00:00 2001 From: SeanOMik Date: Sat, 16 Mar 2024 18:50:22 -0400 Subject: [PATCH 03/15] position the camera in a good position in the scene --- examples/testbed/src/main.rs | 20 +++++++++++++++----- 1 file changed, 15 insertions(+), 5 deletions(-) diff --git a/examples/testbed/src/main.rs b/examples/testbed/src/main.rs index dba017e..0f9d152 100644 --- a/examples/testbed/src/main.rs +++ b/examples/testbed/src/main.rs @@ -1,6 +1,6 @@ -use std::ptr::NonNull; +use std::{cell::Ref, 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, 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; @@ -218,7 +218,13 @@ 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(()) @@ -257,11 +263,15 @@ async fn main() { }; let jiggle_plugin = move |game: &mut Game| { - game.world_mut().add_resource(TpsAccumulator(0.0)); + /* 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(spin_system.into_system()); */ + + /* let mut camera = CameraComponent::new_3d(); + camera.transform.translation += math::Vec3::new(0.0, 0.0, 5.5); + world.spawn(( camera, FreeFlyCamera::default() )); */ //sys.with_system(fps_system); //game.with_system("fixed", sys, &[]); -- 2.40.1 From 4ce21d4db02f00cafb64bfeb26c6976d5c28b5e6 Mon Sep 17 00:00:00 2001 From: SeanOMik Date: Sat, 16 Mar 2024 19:12:32 -0400 Subject: [PATCH 04/15] render: dont send the same material to the gpu multiple times, speeding up gpu texture loading --- lyra-game/src/render/light_cull_compute.rs | 4 ++-- lyra-game/src/render/renderer.rs | 25 ++++++++++++++++------ 2 files changed, 20 insertions(+), 9 deletions(-) diff --git a/lyra-game/src/render/light_cull_compute.rs b/lyra-game/src/render/light_cull_compute.rs index dca54c1..4cca981 100644 --- a/lyra-game/src/render/light_cull_compute.rs +++ b/lyra-game/src/render/light_cull_compute.rs @@ -1,10 +1,10 @@ -use std::{borrow::Cow, mem, num::NonZeroU32, ptr::NonNull, rc::Rc}; +use std::{borrow::Cow, mem, ptr::NonNull, rc::Rc}; use glam::UVec2; use wgpu::{util::DeviceExt, ComputePipeline}; use winit::dpi::PhysicalSize; -use super::{light::LightUniformBuffers, render_buffer::BufferWrapper, renderer::RenderPass, texture::RenderTexture}; +use super::{light::LightUniformBuffers, render_buffer::BufferWrapper, texture::RenderTexture}; struct LightIndicesGridBuffer { indices_buffer: wgpu::Buffer, diff --git a/lyra-game/src/render/renderer.rs b/lyra-game/src/render/renderer.rs index 47eda68..74db284 100755 --- a/lyra-game/src/render/renderer.rs +++ b/lyra-game/src/render/renderer.rs @@ -1,4 +1,5 @@ use std::collections::{HashMap, VecDeque, HashSet}; +use std::ops::Deref; use std::rc::Rc; use std::sync::Arc; use std::borrow::Cow; @@ -10,6 +11,7 @@ use lyra_ecs::query::filter::{Has, Or}; use lyra_ecs::{Entity, Tick}; use lyra_ecs::query::{Entities, TickOf}; use lyra_ecs::World; +use lyra_reflect::resource; use lyra_resource::gltf::GltfScene; use tracing::{debug, warn}; use uuid::Uuid; @@ -20,6 +22,7 @@ use winit::window::Window; use crate::math::Transform; use crate::render::material::MaterialUniform; use crate::render::render_buffer::BufferWrapperBuilder; +use crate::resources; use crate::scene::CameraComponent; use super::camera::{RenderCamera, CameraUniform}; @@ -59,7 +62,7 @@ struct MeshBufferStorage { //#[allow(dead_code)] //render_texture: Option, - material: Option, + material: Option>, // The index of the transform for this entity. // The tuple is structured like this: (transform index, index of transform inside the buffer) @@ -88,6 +91,7 @@ pub struct BasicRenderer { pub render_jobs: VecDeque, mesh_buffers: HashMap, // TODO: clean up left over buffers from deleted entities/components + material_buffers: HashMap>, entity_meshes: HashMap, entity_last_transforms: HashMap, @@ -223,6 +227,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, @@ -347,18 +352,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()), } } -- 2.40.1 From c73c1a7f43e3093b883fd3a24a0a78c619d82c56 Mon Sep 17 00:00:00 2001 From: SeanOMik Date: Sat, 16 Mar 2024 22:58:38 -0400 Subject: [PATCH 05/15] render: fix segfault in LightCullCompute --- examples/testbed/testbed-renderdoc.cap | 4 +- lyra-game/src/game.rs | 1 + lyra-game/src/render/light/mod.rs | 12 ++-- lyra-game/src/render/light_cull_compute.rs | 56 +++++-------------- lyra-game/src/render/material.rs | 8 +-- lyra-game/src/render/render_buffer.rs | 16 +++--- lyra-game/src/render/renderer.rs | 19 ++++--- .../src/render/shaders/light_cull.comp.wgsl | 5 +- lyra-game/src/render/texture.rs | 10 ++-- 9 files changed, 51 insertions(+), 80 deletions(-) diff --git a/examples/testbed/testbed-renderdoc.cap b/examples/testbed/testbed-renderdoc.cap index 7b81a15..cd8a585 100644 --- a/examples/testbed/testbed-renderdoc.cap +++ b/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" } } diff --git a/lyra-game/src/game.rs b/lyra-game/src/game.rs index fe9f9d9..32f7cbd 100755 --- a/lyra-game/src/game.rs +++ b/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(); diff --git a/lyra-game/src/render/light/mod.rs b/lyra-game/src/render/light/mod.rs index 9cdc09b..1f5cfc9 100644 --- a/lyra-game/src/render/light/mod.rs +++ b/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 LightBuffer { 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, 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, } diff --git a/lyra-game/src/render/light_cull_compute.rs b/lyra-game/src/render/light_cull_compute.rs index 4cca981..c928751 100644 --- a/lyra-game/src/render/light_cull_compute.rs +++ b/lyra-game/src/render/light_cull_compute.rs @@ -1,44 +1,28 @@ -use std::{borrow::Cow, mem, 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, texture::RenderTexture}; +use super::{light::LightUniformBuffers, render_buffer::{BindGroupPair, BufferWrapper}, texture::RenderTexture}; struct LightIndicesGridBuffer { indices_buffer: wgpu::Buffer, grid_texture: wgpu::Texture, grid_texture_view: wgpu::TextureView, - bind_group_layout: wgpu::BindGroupLayout, - bind_group: wgpu::BindGroup, + bg_pair: BindGroupPair, } pub(crate) struct LightCullCompute { device: Rc, queue: Rc, pipeline: ComputePipeline, - lights: NonNull, - camera: NonNull, light_indices_grid: LightIndicesGridBuffer, screen_size_buffer: BufferWrapper, - depth_tex: NonNull, } impl LightCullCompute { fn create_grid(device: &wgpu::Device, screen_size: PhysicalSize) -> 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::() * 16 * 16) as u64, - mapped_at_creation: false, - } - ); */ - let light_indices_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor { label: Some("B_LightIndices"), contents: &[0; mem::size_of::() * 16 * 16], @@ -128,12 +112,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, queue: Rc, screen_size: PhysicalSize, lights: &LightUniformBuffers, camera_buffers: &BufferWrapper, depth_texture: &mut RenderTexture) -> Self { + pub fn new(device: Rc, queue: Rc, screen_size: PhysicalSize, 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") @@ -150,16 +133,16 @@ impl LightCullCompute { let light_grid = Self::create_grid(&device, screen_size); + let depth_tex_pair = depth_texture.create_bind_group(&device); let layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { label: Some("PipeLay_LightCull"), bind_group_layouts: &[ - //&depth_texture.bindgroup_pair.as_ref().unwrap().layout, &depth_tex_pair.layout, - camera_buffers.bindgroup_layout().unwrap(), - &lights.bindgroup_layout, - &light_grid.bind_group_layout, + &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,11 +159,8 @@ 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), } } @@ -189,9 +169,7 @@ impl LightCullCompute { &[UVec2::new(size.width, size.height)]); } - pub fn compute(&mut self) { - //self.queue.write_buffer(&self.light_indices_grid.indices_buffer, 0, &[0; mem::size_of::() * 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"), }); @@ -202,17 +180,11 @@ 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); diff --git a/lyra-game/src/render/material.rs b/lyra-game/src/render/material.rs index aa452d2..b505ffc 100755 --- a/lyra-game/src/render/material.rs +++ b/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, } -fn texture_to_render(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: &Arc, i: &Option>) -> Option { +fn texture_to_render(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: &Rc, i: &Option>) -> Option { 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, value: &lyra_resource::gltf::Specular) -> Self { + pub fn from_resource(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Rc, 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, value: &lyra_resource::gltf::Material) -> Self { + pub fn from_resource(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Rc, 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)); diff --git a/lyra-game/src/render/render_buffer.rs b/lyra-game/src/render/render_buffer.rs index df174e0..794e09e 100755 --- a/lyra-game/src/render/render_buffer.rs +++ b/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, + pub layout: Rc, } impl BindGroupPair { - pub fn new_from_layout(device: &wgpu::Device, layout: Arc, entries: &[wgpu::BindGroupEntry<'_>]) -> Self { + pub fn create_bind_group(device: &wgpu::Device, layout: Rc, 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), } } }; diff --git a/lyra-game/src/render/renderer.rs b/lyra-game/src/render/renderer.rs index 74db284..67172f3 100755 --- a/lyra-game/src/render/renderer.rs +++ b/lyra-game/src/render/renderer.rs @@ -1,5 +1,4 @@ use std::collections::{HashMap, VecDeque, HashSet}; -use std::ops::Deref; use std::rc::Rc; use std::sync::Arc; use std::borrow::Cow; @@ -11,7 +10,6 @@ use lyra_ecs::query::filter::{Has, Or}; use lyra_ecs::{Entity, Tick}; use lyra_ecs::query::{Entities, TickOf}; use lyra_ecs::World; -use lyra_reflect::resource; use lyra_resource::gltf::GltfScene; use tracing::{debug, warn}; use uuid::Uuid; @@ -22,7 +20,6 @@ use winit::window::Window; use crate::math::Transform; use crate::render::material::MaterialUniform; use crate::render::render_buffer::BufferWrapperBuilder; -use crate::resources; use crate::scene::CameraComponent; use super::camera::{RenderCamera, CameraUniform}; @@ -103,7 +100,7 @@ pub struct BasicRenderer { camera_buffer: BufferWrapper, //camera_bind_group: wgpu::BindGroup, - bgl_texture: Arc, + bgl_texture: Rc, default_texture: RenderTexture, depth_buffer_texture: RenderTexture, @@ -176,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 { @@ -252,7 +249,7 @@ 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.light_buffers.bind_group_pair.layout, &s.material_buffer.bindgroup_pair.as_ref().unwrap().layout, &s.bgl_texture]))); s.render_pipelines = pipelines; @@ -520,7 +517,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") @@ -582,7 +579,7 @@ 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, &[]); // if this mesh uses indices, use them to draw the mesh @@ -616,7 +613,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); } diff --git a/lyra-game/src/render/shaders/light_cull.comp.wgsl b/lyra-game/src/render/shaders/light_cull.comp.wgsl index 54df3c5..d5dc8d5 100644 --- a/lyra-game/src/render/shaders/light_cull.comp.wgsl +++ b/lyra-game/src/render/shaders/light_cull.comp.wgsl @@ -62,11 +62,8 @@ var u_lights: Lights; @group(3) @binding(0) var u_light_indices: array; -/*@group(3) @binding(1) -var u_light_grid: array>>;*/ - @group(3) @binding(1) -var t_light_grid: texture_storage_2d; // rg32uint = vec2 or vec4(r, g, 0.0, 1.0) +var t_light_grid: texture_storage_2d; @group(4) @binding(0) var u_screen_size: vec2; diff --git a/lyra-game/src/render/texture.rs b/lyra-game/src/render/texture.rs index 154ad11..64e2b59 100755 --- a/lyra-game/src/render/texture.rs +++ b/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, view: &wgpu::TextureView, sampler: &wgpu::Sampler) -> BindGroupPair { + fn create_bind_group_pair(device: &wgpu::Device, layout: Rc, 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, bytes: &[u8], label: &str) -> anyhow::Result { + pub fn from_bytes(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Rc, bytes: &[u8], label: &str) -> anyhow::Result { 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, img: &image::DynamicImage, label: Option<&str>) -> anyhow::Result { + pub fn from_image(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Rc, img: &image::DynamicImage, label: Option<&str>) -> anyhow::Result { 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, texture_res: &ResHandle, label: Option<&str>) -> anyhow::Result { + pub fn from_resource(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Rc, texture_res: &ResHandle, label: Option<&str>) -> anyhow::Result { let texture_ref = texture_res.data_ref().unwrap(); let img = texture_ref.image.data_ref().unwrap(); -- 2.40.1 From 5c1ce809ff0f24502e97961323161484b0b6ac04 Mon Sep 17 00:00:00 2001 From: SeanOMik Date: Sun, 17 Mar 2024 15:20:17 -0400 Subject: [PATCH 06/15] 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 | 4 +- examples/testbed/src/main.rs | 47 ++++++----- lyra-game/src/render/light_cull_compute.rs | 25 +++--- lyra-game/src/render/renderer.rs | 6 +- lyra-game/src/render/shaders/base.wgsl | 78 ++++++++++++++----- .../src/render/shaders/light_cull.comp.wgsl | 68 ++++++++-------- 6 files changed, 143 insertions(+), 85 deletions(-) diff --git a/examples/testbed/src/free_fly_camera.rs b/examples/testbed/src/free_fly_camera.rs index f2fd972..b6f404b 100644 --- a/examples/testbed/src/free_fly_camera.rs +++ b/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, } } diff --git a/examples/testbed/src/main.rs b/examples/testbed/src/main.rs index 0f9d152..3877252 100644 --- a/examples/testbed/src/main.rs +++ b/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::("assets/happy-tree.png").unwrap(); //let antique_camera_model = resman.request::("assets/AntiqueCamera.glb").unwrap(); //let cube_model = resman.request::("assets/cube-texture-bin.glb").unwrap(); - /* let cube_gltf = resman.request::("assets/texture-sep/texture-sep.gltf").unwrap(); - let crate_gltf = resman.request::("assets/crate/crate.gltf").unwrap(); + let cube_gltf = resman.request::("assets/texture-sep/texture-sep.gltf").unwrap(); + /*let crate_gltf = resman.request::("assets/crate/crate.gltf").unwrap(); let separate_gltf = resman.request::("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, - - ambient: 0.0, - diffuse: 7.0, - specular: 1.0, + linear: 0.09, + quadratic: 0.032, + + 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(( diff --git a/lyra-game/src/render/light_cull_compute.rs b/lyra-game/src/render/light_cull_compute.rs index c928751..441ca16 100644 --- a/lyra-game/src/render/light_cull_compute.rs +++ b/lyra-game/src/render/light_cull_compute.rs @@ -6,26 +6,31 @@ use winit::dpi::PhysicalSize; 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, - bg_pair: BindGroupPair, + pub bg_pair: BindGroupPair, } pub(crate) struct LightCullCompute { device: Rc, queue: Rc, pipeline: ComputePipeline, - light_indices_grid: LightIndicesGridBuffer, + pub light_indices_grid: LightIndicesGridBuffer, screen_size_buffer: BufferWrapper, + workgroup_size: glam::UVec2, } impl LightCullCompute { - fn create_grid(device: &wgpu::Device, screen_size: PhysicalSize) -> LightIndicesGridBuffer { + fn create_grid(device: &wgpu::Device, screen_size: PhysicalSize, workgroup_size: glam::UVec2) -> LightIndicesGridBuffer { + let mut contents = Vec::::new(); + let contents_len = workgroup_size.x * workgroup_size.y * mem::size_of::() 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::() * 16 * 16], + contents: &contents, usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, }); @@ -131,8 +136,8 @@ 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); @@ -161,12 +166,14 @@ impl LightCullCompute { pipeline, light_indices_grid: light_grid, screen_size_buffer, + workgroup_size, } } - pub fn update_screen_size(&self, size: PhysicalSize) { + pub fn update_screen_size(&mut self, size: PhysicalSize) { 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) { @@ -187,7 +194,7 @@ impl LightCullCompute { 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); diff --git a/lyra-game/src/render/renderer.rs b/lyra-game/src/render/renderer.rs index 67172f3..819801b 100755 --- a/lyra-game/src/render/renderer.rs +++ b/lyra-game/src/render/renderer.rs @@ -250,7 +250,9 @@ impl BasicRenderer { 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.bgl_texture]))); + &s.bgl_texture, + &s.light_cull_compute.light_indices_grid.bg_pair.layout, + ]))); s.render_pipelines = pipelines; s @@ -582,6 +584,8 @@ impl Renderer for BasicRenderer { 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; diff --git a/lyra-game/src/render/shaders/base.wgsl b/lyra-game/src/render/shaders/base.wgsl index c3c8bc1..7aa3024 100755 --- a/lyra-game/src/render/shaders/base.wgsl +++ b/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, @@ -97,41 +97,79 @@ var t_specular: texture_2d; @group(5) @binding(1) var s_specular: sampler; +@group(6) @binding(0) +var u_light_indices: array; +@group(6) @binding(1) +var t_light_grid: texture_storage_2d; // vec2 + @fragment fn fs_main(in: VertexOutput) -> @location(0) vec4 { - let object_color: vec4 = textureSample(t_diffuse, s_diffuse, in.tex_coords); + /*let object_color: vec4 = textureSample(t_diffuse, s_diffuse, in.tex_coords); let specular_color: vec3 = 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(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(light_object_res, object_color.a);*/ + + let tile_index = vec2(floor(in.clip_position.xy / 16.0)); + let tile: vec2 = textureLoad(t_light_grid, tile_index).xy; + + let object_color: vec4 = textureSample(t_diffuse, s_diffuse, in.tex_coords); + let specular_color: vec3 = textureSample(t_specular, s_specular, in.tex_coords).xyz; + var light_res = vec3(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(light_object_res, object_color.a); + + //return debug_grid(in); +} + +fn debug_grid(in: VertexOutput) -> vec4 { + let tile_index_float: vec2 = in.clip_position.xy / 16.0; + let tile_index = vec2(floor(tile_index_float)); + let tile: vec2 = 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(0.0, 0.0, 0.0, 1.0); + } else { + return vec4(f32(tile_index.x) / 50.0, f32(tile_index.y) / 38.0, 0.0, 1.0); + } } fn blinn_phong_dir_light(world_pos: vec3, world_norm: vec3, dir_light: Light, material: Material, specular_factor: vec3) -> vec3 { diff --git a/lyra-game/src/render/shaders/light_cull.comp.wgsl b/lyra-game/src/render/shaders/light_cull.comp.wgsl index d5dc8d5..4ee4353 100644 --- a/lyra-game/src/render/shaders/light_cull.comp.wgsl +++ b/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 wg_max_depth: atomic; var wg_frustum_planes: array, 6>; // index list of visible light sources for this tile -var wg_visible_light_indices: array; +var wg_visible_light_indices: array; var wg_visible_light_count: atomic; //var view_projection: mat4x4; @group(0) @binding(0) -var t_depthmap: texture_2d; +var t_depthmap: texture_depth_2d; @group(0) @binding(1) var s_depthmap: sampler; @@ -93,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(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(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(depth_float); @@ -143,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 @@ -157,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; - } + }*/ } } } @@ -179,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(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(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 @@ -221,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, 6>, sphere_origin: vec3, radius: f32) -> bool { // to be able to index this array with a non-const value, // it must be defined as a var @@ -241,5 +243,7 @@ fn sphere_inside_frustrum(frustum: array, 6>, sphere_origin: vec3 /// 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, radius: f32, plane: vec4) -> bool { - return dot(plane.xyz, sphere_origin) - plane.w < -radius; + //return dot(plane.xyz, sphere_origin) - plane.w < -radius; + + return dot(vec4(sphere_origin, 0.0), plane) + radius > 0.0; } \ No newline at end of file -- 2.40.1 From 76ec9606ecb8935eb48adbbf991bb38a975c67e8 Mon Sep 17 00:00:00 2001 From: SeanOMik Date: Sun, 17 Mar 2024 16:07:24 -0400 Subject: [PATCH 07/15] render: add some fields to the camera uniform --- lyra-game/src/render/camera.rs | 39 ++++++++++++++----- lyra-game/src/render/renderer.rs | 6 +-- lyra-game/src/render/shaders/base.wgsl | 13 ++++--- .../src/render/shaders/light_cull.comp.wgsl | 24 ++++++------ 4 files changed, 52 insertions(+), 30 deletions(-) diff --git a/lyra-game/src/render/camera.rs b/lyra-game/src/render/camera.rs index 00d904d..882f7c5 100755 --- a/lyra-game/src/render/camera.rs +++ b/lyra-game/src/render/camera.rs @@ -39,9 +39,11 @@ impl Projection { #[derive(Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)] pub struct CameraUniform { /// The view matrix of the camera - pub view_mat: glam::Mat4, + pub view: glam::Mat4, + /// The inverse of the projection matrix of the camera + pub inverse_projection: glam::Mat4, /// The view projection matrix - pub view_proj: glam::Mat4, + pub view_projection: glam::Mat4, /// The position of the camera pub position: glam::Vec3, _padding: u32, @@ -51,8 +53,9 @@ pub struct CameraUniform { impl Default for CameraUniform { fn default() -> Self { Self { - view_mat: glam::Mat4::IDENTITY, - view_proj: glam::Mat4::IDENTITY, + view: glam::Mat4::IDENTITY, + inverse_projection: glam::Mat4::IDENTITY, + view_projection: glam::Mat4::IDENTITY, position: Default::default(), _padding: 0, } @@ -60,10 +63,11 @@ impl Default for CameraUniform { } impl CameraUniform { - pub fn new(view_mat: glam::Mat4, view_proj: glam::Mat4, position: glam::Vec3) -> Self { + pub fn new(view: glam::Mat4, inverse_projection: glam::Mat4, view_projection: glam::Mat4, position: glam::Vec3) -> Self { Self { - view_mat, - view_proj, + view, + inverse_projection, + view_projection, position, _padding: 0 } @@ -101,7 +105,7 @@ impl RenderCamera { /// /// 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) -> (&glam::Mat4, glam::Mat4) { + 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(); @@ -117,7 +121,15 @@ impl RenderCamera { 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, view) + //(&self.view_proj, view) + + CameraUniform { + view, + inverse_projection: proj.inverse(), + view_projection: self.view_proj, + position, + _padding: 0, + } }, CameraProjectionMode::Orthographic => { let target = camera.transform.rotation * glam::Vec3::new(0.0, 0.0, -1.0); @@ -132,7 +144,14 @@ 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, view) + + CameraUniform { + view, + inverse_projection: proj.inverse(), + view_projection: self.view_proj, + position, + _padding: 0, + } }, } } diff --git a/lyra-game/src/render/renderer.rs b/lyra-game/src/render/renderer.rs index 819801b..c960081 100755 --- a/lyra-game/src/render/renderer.rs +++ b/lyra-game/src/render/renderer.rs @@ -504,9 +504,9 @@ impl Renderer for BasicRenderer { } if let Some(camera) = main_world.view_iter::<&mut CameraComponent>().next() { - let (view_proj, view_mat) = self.inuse_camera.calc_view_projection(&camera); - let pos = camera.transform.translation; - let uniform = CameraUniform::new(view_mat, *view_proj, pos); + let uniform = self.inuse_camera.calc_view_projection(&camera); + //let pos = camera.transform.translation; + //let uniform = CameraUniform::new(view_mat, *view_proj, pos); self.camera_buffer.write_buffer(&self.queue, 0, &[uniform]); } else { warn!("Missing camera!"); diff --git a/lyra-game/src/render/shaders/base.wgsl b/lyra-game/src/render/shaders/base.wgsl index 7aa3024..f025ca4 100755 --- a/lyra-game/src/render/shaders/base.wgsl +++ b/lyra-game/src/render/shaders/base.wgsl @@ -20,9 +20,10 @@ struct VertexOutput { } struct CameraUniform { - view_mat: mat4x4, - view_proj: mat4x4, - view_pos: vec3, + view: mat4x4, + projection: mat4x4, + view_projection: mat4x4, + position: vec3, }; struct Light { @@ -63,7 +64,7 @@ fn vs_main( var out: VertexOutput; out.tex_coords = model.tex_coords; - out.clip_position = u_camera.view_proj * u_model_transform * vec4(model.position, 1.0); + out.clip_position = u_camera.view_projection * u_model_transform * vec4(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); @@ -174,7 +175,7 @@ fn debug_grid(in: VertexOutput) -> vec4 { fn blinn_phong_dir_light(world_pos: vec3, world_norm: vec3, dir_light: Light, material: Material, specular_factor: vec3) -> vec3 { 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; @@ -204,7 +205,7 @@ fn blinn_phong_dir_light(world_pos: vec3, world_norm: vec3, dir_light: fn blinn_phong_point_light(world_pos: vec3, world_norm: vec3, point_light: Light, material: Material, specular_factor: vec3) -> vec3 { 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; diff --git a/lyra-game/src/render/shaders/light_cull.comp.wgsl b/lyra-game/src/render/shaders/light_cull.comp.wgsl index 4ee4353..e26f170 100644 --- a/lyra-game/src/render/shaders/light_cull.comp.wgsl +++ b/lyra-game/src/render/shaders/light_cull.comp.wgsl @@ -14,9 +14,10 @@ const LIGHT_TY_SPOT = 2u; // local_invocation_index struct CameraUniform { - view_mat: mat4x4, - view_proj: mat4x4, - view_pos: vec3, + view: mat4x4, + projection: mat4x4, + view_projection: mat4x4, + position: vec3, }; struct Light { @@ -103,11 +104,11 @@ fn cs_main( atomicMin(&wg_min_depth, depth_uint); atomicMax(&wg_max_depth, depth_uint); + workgroupBarrier(); + // convert them back into floats var min_depth: f32 = bitcast(wg_min_depth); var max_depth: f32 = bitcast(wg_max_depth); - - workgroupBarrier(); // Create the frustum planes that will be used for this time if (local_invocation_index == 0u) { @@ -124,14 +125,14 @@ fn cs_main( // convert the side and top planes from clip to view space for (var i = 0u; i < 4u; i++) { - wg_frustum_planes[i] *= u_camera.view_proj; + wg_frustum_planes[i] *= u_camera.view_projection; wg_frustum_planes[i] /= length(wg_frustum_planes[i].xyz); } // convert near and far planes from clip to view space - wg_frustum_planes[4] *= u_camera.view_mat; + wg_frustum_planes[4] *= u_camera.view; wg_frustum_planes[4] /= length(wg_frustum_planes[4].xyz); - wg_frustum_planes[5] *= u_camera.view_mat; + wg_frustum_planes[5] *= u_camera.view; wg_frustum_planes[5] /= length(wg_frustum_planes[5].xyz); } @@ -148,7 +149,7 @@ fn cs_main( // find the light index to check on this thread, make sure we're not trying to test // for more lights than we have. var light_index = i * thread_count + local_invocation_index; - if (light_index > u_lights.light_count) { + if (light_index >= u_lights.light_count) { break; } @@ -157,14 +158,15 @@ fn cs_main( var radius = light.range; if (light.light_ty == LIGHT_TY_DIRECTIONAL) { - add_light(light_index); + //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 + add_light(light_index); + // TODO: spotlights if (!sphere_inside_plane(position, radius, wg_frustum_planes[4])) { - add_light(light_index); /*var offset: u32 = wg_visible_light_count; if (offset < MAX_TILE_VISIBLE_LIGHTS) { -- 2.40.1 From 014abcf7e60c9c4a10c582183a27a39c3af38f73 Mon Sep 17 00:00:00 2001 From: SeanOMik Date: Sun, 17 Mar 2024 18:11:38 -0400 Subject: [PATCH 08/15] render: fix the tile frustum used for culling lights --- examples/testbed/src/main.rs | 5 +- lyra-game/src/render/light/mod.rs | 2 +- lyra-game/src/render/light_cull_compute.rs | 38 +++- lyra-game/src/render/shaders/base.wgsl | 4 +- .../src/render/shaders/light_cull.comp.wgsl | 169 +++++++++--------- shell.nix | 1 + 6 files changed, 131 insertions(+), 88 deletions(-) diff --git a/examples/testbed/src/main.rs b/examples/testbed/src/main.rs index 3877252..f7e8e54 100644 --- a/examples/testbed/src/main.rs +++ b/examples/testbed/src/main.rs @@ -170,7 +170,10 @@ async fn main() { specular: 1.3, }, Transform::new( - Vec3::new(-5.0, 1.0, -1.28), + //Vec3::new(-5.0, 1.0, -1.28), + Vec3::new(-5.0, 1.0, -0.28), + //Vec3::new(-10.0, 0.94, -0.28), + Quat::IDENTITY, Vec3::new(0.25, 0.25, 0.25), ), diff --git a/lyra-game/src/render/light/mod.rs b/lyra-game/src/render/light/mod.rs index 1f5cfc9..eaebbaa 100644 --- a/lyra-game/src/render/light/mod.rs +++ b/lyra-game/src/render/light/mod.rs @@ -247,7 +247,7 @@ impl LightUniform { _padding3: 0, color: light.color, - range: 2.0, + range: 1.5, intensity: 1.0, spot_cutoff: 0.0, diff --git a/lyra-game/src/render/light_cull_compute.rs b/lyra-game/src/render/light_cull_compute.rs index 441ca16..089c313 100644 --- a/lyra-game/src/render/light_cull_compute.rs +++ b/lyra-game/src/render/light_cull_compute.rs @@ -7,6 +7,7 @@ use winit::dpi::PhysicalSize; use super::{light::LightUniformBuffers, render_buffer::{BindGroupPair, BufferWrapper}, texture::RenderTexture}; pub(crate) struct LightIndicesGridBuffer { + index_counter_buffer: wgpu::Buffer, indices_buffer: wgpu::Buffer, grid_texture: wgpu::Texture, grid_texture_view: wgpu::TextureView, @@ -34,6 +35,12 @@ impl LightCullCompute { 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 { @@ -57,15 +64,27 @@ impl LightCullCompute { 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"), }); - // TODO: shrink the texture to match the amount of grid cells that the shader actually uses + // TODO: resize texture when screen is resized let size = wgpu::Extent3d { - width: screen_size.width, - height: screen_size.height, + width: workgroup_size.x, + height: workgroup_size.y, depth_or_array_layers: 1, }; let grid_texture = device.create_texture( @@ -108,12 +127,23 @@ impl LightCullCompute { 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, diff --git a/lyra-game/src/render/shaders/base.wgsl b/lyra-game/src/render/shaders/base.wgsl index f025ca4..03c990b 100755 --- a/lyra-game/src/render/shaders/base.wgsl +++ b/lyra-game/src/render/shaders/base.wgsl @@ -21,7 +21,7 @@ struct VertexOutput { struct CameraUniform { view: mat4x4, - projection: mat4x4, + inverse_projection: mat4x4, view_projection: mat4x4, position: vec3, }; @@ -126,6 +126,8 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4 { return vec4(light_object_res, object_color.a);*/ + + let tile_index = vec2(floor(in.clip_position.xy / 16.0)); let tile: vec2 = textureLoad(t_light_grid, tile_index).xy; diff --git a/lyra-game/src/render/shaders/light_cull.comp.wgsl b/lyra-game/src/render/shaders/light_cull.comp.wgsl index e26f170..e45f27a 100644 --- a/lyra-game/src/render/shaders/light_cull.comp.wgsl +++ b/lyra-game/src/render/shaders/light_cull.comp.wgsl @@ -1,21 +1,14 @@ -const BLOCK_SIZE: i32 = 16; +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; -// Possible computer shader inputs: -// -// local_invocation_id -// workgroup_id -// global_invocation_id -// num_workgroups -// local_invocation_index - struct CameraUniform { view: mat4x4, - projection: mat4x4, + inverse_projection: mat4x4, + //projection: mat4x4, view_projection: mat4x4, position: vec3, }; @@ -42,14 +35,13 @@ struct Lights { var wg_min_depth: atomic; var wg_max_depth: atomic; +var wg_light_index_start: atomic; var wg_frustum_planes: array, 6>; // index list of visible light sources for this tile var wg_visible_light_indices: array; var wg_visible_light_count: atomic; -//var view_projection: mat4x4; - @group(0) @binding(0) var t_depthmap: texture_depth_2d; @group(0) @binding(1) @@ -65,6 +57,8 @@ var u_lights: Lights; var u_light_indices: array; @group(3) @binding(1) var t_light_grid: texture_storage_2d; +@group(3) @binding(2) +var u_light_index_counter: atomic; @group(4) @binding(0) var u_screen_size: vec2; @@ -78,12 +72,6 @@ fn cs_main( @builtin(num_workgroups) num_workgroups: vec3, @builtin(local_invocation_index) local_invocation_index: u32, ) { - //var location = vec2(global_invocation_id.xy); - var item_id = vec2(local_invocation_id.xy); - var tile_id = vec2(workgroup_id.xy); - var tile_number = vec2(num_workgroups.xy); - var index = tile_id.y * tile_number.x + tile_id.x; - // Initialize some shared global values for depth and light count if (local_invocation_index == 0u) { wg_min_depth = 0xFFFFFFFu; @@ -112,28 +100,41 @@ fn cs_main( // Create the frustum planes that will be used for this time if (local_invocation_index == 0u) { - var negative_step = (2.0 * vec2(tile_id)) / vec2(tile_number); - var positive_step = (2.0 * vec2(tile_id) + vec2(1.0, 1.0)) / vec2(tile_number); + // Compute the 4 corner points on the far clipping plane to use as the frustum vertices. + var screen_space: array, 4>; - // z in the vec4 is the distance from the center of the tile - wg_frustum_planes[0] = vec4(1.0, 0.0, 0.0, 1.0 - negative_step.x); // left - wg_frustum_planes[1] = vec4(-1.0, 0.0, 0.0, -1.0 + positive_step.x); // right - wg_frustum_planes[2] = vec4(0.0, -1.0, 0.0, 1.0 - negative_step.y); // bottom - wg_frustum_planes[3] = vec4(0.0, -1.0, 0.0, -1.0 + positive_step.y); // top - wg_frustum_planes[4] = vec4(0.0, 0.0, -1.0, -min_depth); // near plane - wg_frustum_planes[5] = vec4(0.0, 0.0, 1.0, max_depth); // far plane + // top left point + var temp: vec2 = workgroup_id.xy * BLOCK_SIZE; + screen_space[0] = vec4(f32(temp.x), f32(temp.y), -1.0, 1.0); - // convert the side and top planes from clip to view space + // top right point + var temp2 = vec2(f32(workgroup_id.x) + 1.0, f32(workgroup_id.y)) * f32(BLOCK_SIZE); + screen_space[1] = vec4(temp2.x, temp2.y, -1.0, 1.0); + + // bottom left point + temp2 = vec2(f32(workgroup_id.x), f32(workgroup_id.y) + 1.0) * f32(BLOCK_SIZE); + screen_space[2] = vec4(temp2.x, temp2.y, -1.0, 1.0); + + // bottom right point + temp2 = vec2(f32(workgroup_id.x) + 1.0, f32(workgroup_id.y) + 1.0) * f32(BLOCK_SIZE); + screen_space[3] = vec4(temp2.x, temp2.y, -1.0, 1.0); + + // convert screenspace to view space + var view_space: array, 4>; for (var i = 0u; i < 4u; i++) { - wg_frustum_planes[i] *= u_camera.view_projection; - wg_frustum_planes[i] /= length(wg_frustum_planes[i].xyz); + view_space[i] = screen_to_view(screen_space[i]).xyz; } - // convert near and far planes from clip to view space - wg_frustum_planes[4] *= u_camera.view; - wg_frustum_planes[4] /= length(wg_frustum_planes[4].xyz); - wg_frustum_planes[5] *= u_camera.view; - wg_frustum_planes[5] /= length(wg_frustum_planes[5].xyz); + // View space eye is always at the origin + let eye_pos = vec3(0.0, 0.0, 0.0); + + wg_frustum_planes[0] = compute_plane(eye_pos, view_space[2], view_space[0]); // left plane + wg_frustum_planes[1] = compute_plane(eye_pos, view_space[1], view_space[3]); // right plane + wg_frustum_planes[2] = compute_plane(eye_pos, view_space[0], view_space[1]); // top plane + wg_frustum_planes[3] = compute_plane(eye_pos, view_space[3], view_space[2]); // bottom plane + + wg_frustum_planes[4] = vec4(0.0, 0.0, -1.0, -min_depth); + wg_frustum_planes[5] = vec4(0.0, 0.0, 1.0, -max_depth); } workgroupBarrier(); @@ -143,38 +144,27 @@ 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 pass_count = (u_lights.light_count + thread_count - 1u) / thread_count; - for (var i = 0u; i < pass_count; i++) { - // find the light index to check on this thread, make sure we're not trying to test - // for more lights than we have. - var light_index = i * thread_count + local_invocation_index; - if (light_index >= u_lights.light_count) { - break; - } + for (var i = local_invocation_index; i < u_lights.light_count; i += BLOCK_SIZE * BLOCK_SIZE) { + let light_index = i; - var light = u_lights.data[light_index]; - var position = light.position; - var radius = light.range; + let light = u_lights.data[light_index]; + let position_vec4 = u_camera.view * vec4(light.position, 1.0); + let position = position_vec4.xyz; + let radius = light.range; if (light.light_ty == LIGHT_TY_DIRECTIONAL) { - //add_light(light_index); + 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 add_light(light_index); - // TODO: spotlights if (!sphere_inside_plane(position, radius, wg_frustum_planes[4])) { - /*var offset: u32 = wg_visible_light_count; - - if (offset < MAX_TILE_VISIBLE_LIGHTS) { - atomicAdd(&wg_visible_light_count, 1u); - wg_visible_light_indices[offset] = light_index; - }*/ + } } + // TODO: spotlights } workgroupBarrier(); @@ -183,41 +173,27 @@ 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 - textureStore(t_light_grid, workgroup_id.xy, vec4(offset, wg_visible_light_count, 0u, 1u)); + wg_light_index_start = atomicAdd(&u_light_index_counter, wg_visible_light_count); + textureStore(t_light_grid, workgroup_id.xy, vec4(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. - var indices_offset = u32(index) * MAX_TILE_VISIBLE_LIGHTS; - //var pass_count = (wg_visible_light_count + thread_count - 1) / thread_count; - for (var i = 0u; i < pass_count; i++) { - // find the light index to check on this thread, make sure we're not trying to test - // for more lights than we have. - //var light_index: u32 = i * thread_count + local_invocation_index; - /*if (light_index > u_lights.light_count) { - u_visible_light_indices - break; - }*/ - - var offset = indices_offset + i; - if (offset >= wg_visible_light_count) { - // stop if we're over the over the amount of lights we saw - break; - } - - u_light_indices[offset] = wg_visible_light_indices[i]; + 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; + //var offset: u32 = wg_visible_light_count; - if (offset < MAX_TILE_VISIBLE_LIGHTS) { - atomicAdd(&wg_visible_light_count, 1u); + 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; } @@ -232,7 +208,7 @@ fn sphere_inside_frustrum(frustum: array, 6>, sphere_origin: vec3 // only check the sides of the frustum for (var i = 0u; i < 4u; i++) { - if (!sphere_inside_plane(sphere_origin, radius, frustum_v[i])) { + if (sphere_inside_plane(sphere_origin, radius, frustum_v[i])) { return false; } } @@ -245,7 +221,38 @@ fn sphere_inside_frustrum(frustum: array, 6>, sphere_origin: vec3 /// 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, radius: f32, plane: vec4) -> bool { - //return dot(plane.xyz, sphere_origin) - plane.w < -radius; + return dot(plane.xyz, sphere_origin) - plane.w < -radius; +} - return dot(vec4(sphere_origin, 0.0), plane) + radius > 0.0; +fn clip_to_view(clip: vec4) -> vec4 { + // view space position + var view = u_camera.inverse_projection * clip; + + // perspective projection + return view / view.w; +} + +fn screen_to_view(screen: vec4) -> vec4 { + // convert to normalized texture coordinates + let tex_coord = screen.xy / vec2(u_screen_size); + + // convert to clip space + let clip = vec4( vec2(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: vec3, p1: vec3, p2: vec3) -> vec4 { + let v0 = p1 - p0; + let v2 = p2 - p0; + + var plane = vec4(normalize(cross(v0, v2)), 0.0); + + // find the distance to the origin + plane.w = dot(plane.xyz, p0); + + return plane; } \ No newline at end of file diff --git a/shell.nix b/shell.nix index 551dc4c..87fd52d 100755 --- a/shell.nix +++ b/shell.nix @@ -13,6 +13,7 @@ mkShell rec { mold udev lua5_4_compat + rustup ]; buildInputs = [ udev alsa-lib libGL gcc -- 2.40.1 From 834a864544f17b77584bb4e8cd55c44ab657fb62 Mon Sep 17 00:00:00 2001 From: SeanOMik Date: Tue, 19 Mar 2024 20:26:15 -0400 Subject: [PATCH 09/15] render: get forward+ rendering working with multiple light sources --- lyra-game/src/render/light/mod.rs | 15 ++------- lyra-game/src/render/light_cull_compute.rs | 12 +++++-- lyra-game/src/render/shaders/base.wgsl | 5 ++- .../src/render/shaders/light_cull.comp.wgsl | 32 ++++++++++--------- 4 files changed, 31 insertions(+), 33 deletions(-) diff --git a/lyra-game/src/render/light/mod.rs b/lyra-game/src/render/light/mod.rs index eaebbaa..f523263 100644 --- a/lyra-game/src/render/light/mod.rs +++ b/lyra-game/src/render/light/mod.rs @@ -215,15 +215,12 @@ pub(crate) enum LightType { #[repr(C)] #[derive(Default, Debug, Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)] pub(crate) struct LightUniform { - pub light_type: u32, // LightType - pub enabled: u32, // bool - pub _padding: [u32; 2], pub position: glam::Vec3, - pub _padding2: u32, + pub light_type: u32, // LightType pub direction: glam::Vec3, - pub _padding3: 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 range: f32, @@ -240,11 +237,8 @@ impl LightUniform { Self { light_type: LightType::Point as u32, enabled: true as u32, // TODO - _padding: [0; 2], position: transform.translation, - _padding2: 0, direction: transform.forward(), - _padding3: 0, color: light.color, range: 1.5, @@ -261,11 +255,8 @@ impl LightUniform { Self { light_type: LightType::Directional as u32, enabled: true as u32, // TODO: take from component - _padding: [0; 2], position: transform.translation, - _padding2: 0, direction: transform.forward(), - _padding3: 0, color: light.color, range: 0.0, diff --git a/lyra-game/src/render/light_cull_compute.rs b/lyra-game/src/render/light_cull_compute.rs index 089c313..f65018d 100644 --- a/lyra-game/src/render/light_cull_compute.rs +++ b/lyra-game/src/render/light_cull_compute.rs @@ -24,9 +24,10 @@ pub(crate) struct LightCullCompute { } impl LightCullCompute { - fn create_grid(device: &wgpu::Device, screen_size: PhysicalSize, workgroup_size: glam::UVec2) -> LightIndicesGridBuffer { + /// Create the LightIndiciesGridBuffer object + fn create_grid(device: &wgpu::Device, workgroup_size: glam::UVec2) -> LightIndicesGridBuffer { let mut contents = Vec::::new(); - let contents_len = workgroup_size.x * workgroup_size.y * mem::size_of::() as u32; + let contents_len = workgroup_size.x * workgroup_size.y * 200 * mem::size_of::() as u32; contents.resize(contents_len as _, 0); let light_indices_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor { @@ -167,7 +168,7 @@ impl LightCullCompute { }); 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 light_grid = Self::create_grid(&device, workgroup_size); let depth_tex_pair = depth_texture.create_bind_group(&device); @@ -228,5 +229,10 @@ impl LightCullCompute { } 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])); } } \ No newline at end of file diff --git a/lyra-game/src/render/shaders/base.wgsl b/lyra-game/src/render/shaders/base.wgsl index 03c990b..a7a3a19 100755 --- a/lyra-game/src/render/shaders/base.wgsl +++ b/lyra-game/src/render/shaders/base.wgsl @@ -27,11 +27,10 @@ struct CameraUniform { }; struct Light { - light_ty: u32, - enabled: u32, - position: vec3, + light_ty: u32, direction: vec3, + enabled: u32, color: vec3, range: f32, diff --git a/lyra-game/src/render/shaders/light_cull.comp.wgsl b/lyra-game/src/render/shaders/light_cull.comp.wgsl index e45f27a..54c94ac 100644 --- a/lyra-game/src/render/shaders/light_cull.comp.wgsl +++ b/lyra-game/src/render/shaders/light_cull.comp.wgsl @@ -14,11 +14,10 @@ struct CameraUniform { }; struct Light { - light_ty: u32, - enabled: u32, - position: vec3, + light_ty: u32, direction: vec3, + enabled: u32, color: vec3, range: f32, @@ -148,23 +147,26 @@ fn cs_main( let light_index = i; let light = u_lights.data[light_index]; - let position_vec4 = u_camera.view * vec4(light.position, 1.0); - let position = position_vec4.xyz; - let radius = light.range; - 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 + if (light.enabled == 1u) { + let position_vec4 = u_camera.view * vec4(light.position, 1.0); + let position = position_vec4.xyz; + let radius = light.range; - add_light(light_index); + 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 - if (!sphere_inside_plane(position, radius, wg_frustum_planes[4])) { - + add_light(light_index); + + if (!sphere_inside_plane(position, radius, wg_frustum_planes[4])) { + + } } + // TODO: spotlights } - // TODO: spotlights } workgroupBarrier(); -- 2.40.1 From f63a7ae86a2df6d73af315017a3a30761393ab2d Mon Sep 17 00:00:00 2001 From: SeanOMik Date: Tue, 19 Mar 2024 21:40:08 -0400 Subject: [PATCH 10/15] include correct rust install in nix-shell --- rust-toolchain.toml | 7 +++--- shell.nix | 52 +++++++++++++++++++++++---------------------- 2 files changed, 30 insertions(+), 29 deletions(-) diff --git a/rust-toolchain.toml b/rust-toolchain.toml index 5b28b38..78cfb42 100644 --- a/rust-toolchain.toml +++ b/rust-toolchain.toml @@ -1,5 +1,4 @@ [toolchain] -channel = "nightly-2023-11-21" -#components = [ "rustfmt", "rustc-dev" ] -targets = [ "x86_64-unknown-linux-gnu" ] -#profile = "minimal" \ No newline at end of file +channel = "nightly" +date = "2023-11-21" +targets = [ "x86_64-unknown-linux-gnu" ] \ No newline at end of file diff --git a/shell.nix b/shell.nix index 87fd52d..d0c65e9 100755 --- a/shell.nix +++ b/shell.nix @@ -1,25 +1,27 @@ -{ pkgs ? import { } }: - -with pkgs; - -mkShell rec { - nativeBuildInputs = [ - pkg-config - openssl - wasm-pack - trunk - valgrind - heaptrack - mold - udev - lua5_4_compat - rustup - ]; - 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; -} \ No newline at end of file +let + moz_overlay = import (builtins.fetchTarball https://github.com/mozilla/nixpkgs-mozilla/archive/master.tar.gz); + nixpkgs = import { 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 + ]; + 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; + } -- 2.40.1 From 65ff7c4f23a782e34c44478cedb7c86d5ae51be6 Mon Sep 17 00:00:00 2001 From: SeanOMik Date: Tue, 19 Mar 2024 22:40:15 -0400 Subject: [PATCH 11/15] render: retrieve light properties from components --- examples/testbed/src/main.rs | 61 +++++++++---------- lyra-game/src/render/light/directional.rs | 19 ++++-- lyra-game/src/render/light/mod.rs | 20 +++--- lyra-game/src/render/light/point.rs | 25 +++++--- lyra-game/src/render/light_cull_compute.rs | 1 + lyra-game/src/render/shaders/base.wgsl | 19 +----- .../src/render/shaders/light_cull.comp.wgsl | 1 + 7 files changed, 72 insertions(+), 74 deletions(-) diff --git a/examples/testbed/src/main.rs b/examples/testbed/src/main.rs index f7e8e54..a84ac19 100644 --- a/examples/testbed/src/main.rs +++ b/examples/testbed/src/main.rs @@ -120,21 +120,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); @@ -142,36 +127,27 @@ async fn main() { light_tran.rotate_y(math::Angle::Degrees(25.0)); world.spawn(( DirectionalLight { - color: Vec3::new(1.0, 1.0, 1.0), - ambient: 0.3, - diffuse: 1.0, - specular: 1.3, + enabled: true, + color: Vec3::ONE, + intensity: 0.35 + //..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(( PointLight { + enabled: true, color: Vec3::new(0.0, 0.0, 1.0), - - intensity: 3.3, - - constant: 1.0, - linear: 0.09, - quadratic: 0.032, - - ambient: 0.2, - diffuse: 1.0, - specular: 1.3, + 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.28), + Vec3::new(-5.0, 1.0, -0.0), //Vec3::new(-10.0, 0.94, -0.28), Quat::IDENTITY, @@ -179,6 +155,25 @@ async fn main() { ), 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(), + )); } /* { diff --git a/lyra-game/src/render/light/directional.rs b/lyra-game/src/render/light/directional.rs index d3360af..8411b60 100644 --- a/lyra-game/src/render/light/directional.rs +++ b/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, -} \ No newline at end of file +impl Default for DirectionalLight { + fn default() -> Self { + Self { + enabled: true, + color: glam::Vec3::new(1.0, 1.0, 1.0), + intensity: 1.0, + } + } +} diff --git a/lyra-game/src/render/light/mod.rs b/lyra-game/src/render/light/mod.rs index f523263..7d94e2e 100644 --- a/lyra-game/src/render/light/mod.rs +++ b/lyra-game/src/render/light/mod.rs @@ -216,7 +216,7 @@ pub(crate) enum LightType { #[derive(Default, Debug, Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)] pub(crate) struct LightUniform { pub position: glam::Vec3, - pub light_type: u32, // LightType + pub light_type: u32, // enum LightType pub direction: glam::Vec3, pub enabled: u32, // bool pub color: glam::Vec3, @@ -225,47 +225,45 @@ pub(crate) struct LightUniform { pub range: f32, pub intensity: f32, + pub smoothness: f32, pub spot_cutoff: f32, pub spot_outer_cutoff: f32, - - pub _padding4: u32, } impl LightUniform { pub fn from_point_light_bundle(light: &PointLight, transform: &Transform) -> Self { Self { light_type: LightType::Point as u32, - enabled: true as u32, // TODO + enabled: light.enabled as u32, position: transform.translation, direction: transform.forward(), color: light.color, - range: 1.5, - intensity: 1.0, + range: light.range, + intensity: light.intensity, + smoothness: light.smoothness, spot_cutoff: 0.0, spot_outer_cutoff: 0.0, - _padding4: 0, } } pub fn from_directional_bundle(light: &DirectionalLight, transform: &Transform) -> Self { Self { light_type: LightType::Directional as u32, - enabled: true as u32, // TODO: take from component + enabled: light.enabled as u32, position: transform.translation, direction: transform.forward(), color: light.color, range: 0.0, - intensity: 0.0, + intensity: light.intensity, + smoothness: 0.0, spot_cutoff: 0.0, spot_outer_cutoff: 0.0, - - _padding4: 0, } } diff --git a/lyra-game/src/render/light/point.rs b/lyra-game/src/render/light/point.rs index 64f7688..4aa43ef 100644 --- a/lyra-game/src/render/light/point.rs +++ b/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 linear: f32, - pub quadratic: f32, - pub ambient: f32, - pub diffuse: f32, - pub specular: f32, -} \ No newline at end of file + pub smoothness: f32, +} + +impl Default for PointLight { + fn default() -> Self { + Self { + enabled: true, + color: glam::Vec3::new(1.0, 1.0, 1.0), + range: 1.0, + intensity: 1.0, + smoothness: 0.75, + } + } +} diff --git a/lyra-game/src/render/light_cull_compute.rs b/lyra-game/src/render/light_cull_compute.rs index f65018d..98b742e 100644 --- a/lyra-game/src/render/light_cull_compute.rs +++ b/lyra-game/src/render/light_cull_compute.rs @@ -6,6 +6,7 @@ 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, diff --git a/lyra-game/src/render/shaders/base.wgsl b/lyra-game/src/render/shaders/base.wgsl index a7a3a19..aa907a7 100755 --- a/lyra-game/src/render/shaders/base.wgsl +++ b/lyra-game/src/render/shaders/base.wgsl @@ -35,6 +35,7 @@ struct Light { range: f32, intensity: f32, + smoothness: f32, spot_cutoff: f32, spot_outer_cutoff: f32, @@ -200,7 +201,7 @@ fn blinn_phong_dir_light(world_pos: vec3, world_norm: vec3, dir_light: diffuse_color *= dir_light.diffuse; 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, world_norm: vec3, point_light: Light, material: Material, specular_factor: vec3) -> vec3 { @@ -226,22 +227,8 @@ fn blinn_phong_point_light(world_pos: vec3, world_norm: vec3, point_li var specular_color = specular_strength * (light_color * specular_factor); //// end of specular //// - // TODO: Point light range let distance = length(light_pos - world_pos); - // TODO: make smoothness in this a configurable value - // 0.75 is the smoothness or falloff - let attenuation = 1.0 - smoothstep(point_light.range * 0.75, point_light.range, distance); - - //// point light attenuation //// - /*let distance = length(light_pos - world_pos); - let attenuation = 1.0 / (point_light.constant + point_light.linear * distance + - point_light.quadratic * (distance * distance)); - - //// end of point light attenuation //// - - ambient_color *= point_light.ambient * attenuation; - diffuse_color *= point_light.diffuse * attenuation; - specular_color *= point_light.specular * attenuation;*/ + let attenuation = 1.0 - smoothstep(point_light.range * point_light.smoothness, point_light.range, distance); ambient_color *= attenuation; diffuse_color *= attenuation; diff --git a/lyra-game/src/render/shaders/light_cull.comp.wgsl b/lyra-game/src/render/shaders/light_cull.comp.wgsl index 54c94ac..d8e481b 100644 --- a/lyra-game/src/render/shaders/light_cull.comp.wgsl +++ b/lyra-game/src/render/shaders/light_cull.comp.wgsl @@ -22,6 +22,7 @@ struct Light { range: f32, intensity: f32, + smoothness: f32, spot_cutoff: f32, spot_outer_cutoff: f32, -- 2.40.1 From f0b413d9aeb93087b7504cb5b4eed76e532918c1 Mon Sep 17 00:00:00 2001 From: SeanOMik Date: Wed, 20 Mar 2024 11:41:40 -0400 Subject: [PATCH 12/15] render: resize light grid with window, improve light buffer, add spot lights to the light cull compute Spot lights are buggy. They get culled when they shouldn't be, maybe still an issue with the light grid :( --- examples/testbed/src/main.rs | 21 +- lyra-game/src/render/light/mod.rs | 285 +++++------------- lyra-game/src/render/light/spotlight.rs | 23 +- lyra-game/src/render/light_cull_compute.rs | 11 +- lyra-game/src/render/shaders/base.wgsl | 47 ++- .../src/render/shaders/light_cull.comp.wgsl | 65 +++- shell.nix | 8 +- 7 files changed, 242 insertions(+), 218 deletions(-) diff --git a/examples/testbed/src/main.rs b/examples/testbed/src/main.rs index a84ac19..eaa931a 100644 --- a/examples/testbed/src/main.rs +++ b/examples/testbed/src/main.rs @@ -127,7 +127,7 @@ async fn main() { light_tran.rotate_y(math::Angle::Degrees(25.0)); world.spawn(( DirectionalLight { - enabled: true, + enabled: false, color: Vec3::ONE, intensity: 0.35 //..Default::default() @@ -174,6 +174,25 @@ async fn main() { ), cube_mesh.clone(), )); + + /* world.spawn(( + SpotLight { + enabled: true, + color: Vec3::new(1.0, 0.0, 0.0), + intensity: 1.0, + range: 1.5, + ..Default::default() + }, + Transform::new( + Vec3::new(0.0, 0.2, -1.5), + //Vec3::new(-5.0, 1.0, -0.28), + //Vec3::new(-10.0, 0.94, -0.28), + + Quat::IDENTITY, + Vec3::new(0.15, 0.15, 0.15), + ), + cube_mesh.clone(), + )); */ } /* { diff --git a/lyra-game/src/render/light/mod.rs b/lyra-game/src/render/light/mod.rs index 7d94e2e..ca37007 100644 --- a/lyra-game/src/render/light/mod.rs +++ b/lyra-game/src/render/light/mod.rs @@ -6,7 +6,7 @@ use lyra_ecs::{Entity, Tick, World, query::{Entities, TickOf}}; pub use point::*; pub use spotlight::*; -use std::{collections::{HashMap, VecDeque}, marker::PhantomData}; +use std::{collections::{HashMap, VecDeque}, marker::PhantomData, mem}; use crate::math::Transform; @@ -103,13 +103,14 @@ pub(crate) struct LightUniformBuffers { pub buffer: wgpu::Buffer, pub bind_group_pair: BindGroupPair, pub light_indexes: HashMap, + dead_indices: VecDeque, pub current_light_idx: u32, } impl LightUniformBuffers { pub fn new(device: &wgpu::Device) -> Self { let limits = device.limits(); - // TODO: check this limit somehow + // TODO: ensure we dont write over this limit let max_buffer_sizes = (limits.max_uniform_buffer_binding_size as u64) / 2; let buffer = device.create_buffer( @@ -161,45 +162,88 @@ impl LightUniformBuffers { bind_group_pair: BindGroupPair::new(bindgroup, bindgroup_layout), light_indexes: Default::default(), current_light_idx: 0, + dead_indices: VecDeque::new(), } } + /// Returns the index for the entity, and if this index is new + fn get_index_for(&mut self, missed: &mut HashMap, entity: Entity) -> (bool, u32) { + let idx = missed.remove(&entity) + .map(|v| (false, v)) + .or_else(|| + self.dead_indices.pop_front() + .map(|v| (true, v)) + ) + .unwrap_or_else(|| { + let t = self.current_light_idx; + self.current_light_idx += 1; + (true, t) + }); + idx + } + pub fn update_lights(&mut self, queue: &wgpu::Queue, world_tick: Tick, world: &World) { - let mut lights = LightsUniform::default(); + // used to detect what lights were removed + let mut missed_lights: HashMap = self.light_indexes.drain().collect(); for (entity, point_light, transform, light_epoch, transform_epoch) in world.view_iter::<(Entities, &PointLight, &Transform, TickOf, TickOf)>() { - // TODO: dont update light every frame - let idx = *self.light_indexes.entry(entity) - .or_insert_with(|| { - let t = self.current_light_idx; - self.current_light_idx += 1; - t - }) as usize; + let (new, idx) = self.get_index_for(&mut missed_lights, entity); + self.light_indexes.insert(entity, idx); - let uniform = LightUniform::from_point_light_bundle(&point_light, &transform); - lights.data[idx] = uniform; + 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::() * 4 + mem::size_of::() * idx as usize; + queue.write_buffer(&self.buffer, offset as _, bytemuck::cast_slice(&[uniform])); + } } - if let Some((entity, dir_light, transform)) = - world.view_iter::<(Entities, &DirectionalLight, &Transform)>().next() { + for (entity, spot_light, transform, light_epoch, transform_epoch) + in world.view_iter::<(Entities, &SpotLight, &Transform, TickOf, TickOf)>() { + + let (new, idx) = self.get_index_for(&mut missed_lights, entity); + self.light_indexes.insert(entity, idx); + + if new || light_epoch == world_tick || transform_epoch == world_tick { + let uniform = LightUniform::from_spot_light_bundle(&spot_light, &transform); + + let offset = mem::size_of::() * 4 + mem::size_of::() * idx as usize; + queue.write_buffer(&self.buffer, offset as _, bytemuck::cast_slice(&[uniform])); + } + } + + for (entity, dir_light, transform, light_epoch, transform_epoch) + in world.view_iter::<(Entities, &DirectionalLight, &Transform, TickOf, TickOf)>() { - let idx = *self.light_indexes.entry(entity) - .or_insert_with(|| { - let t = self.current_light_idx; - self.current_light_idx += 1; - t - }) as usize; - - let uniform = LightUniform::from_directional_bundle(&dir_light, &transform); - lights.data[idx] = uniform; + let (new, idx) = self.get_index_for(&mut missed_lights, entity); + self.light_indexes.insert(entity, idx); + + if new || light_epoch == world_tick || transform_epoch == world_tick { + let uniform = LightUniform::from_directional_bundle(&dir_light, &transform); + let offset = mem::size_of::() * 4 + mem::size_of::() * idx as usize; + queue.write_buffer(&self.buffer, offset as _, bytemuck::cast_slice(&[uniform])); + } } - lights.light_count = self.light_indexes.len() as u32; + // anything left in missed_lights were lights that were deleted + let len = missed_lights.len(); + self.dead_indices.reserve(len); - // update the light count in the struct - queue.write_buffer(&self.buffer, 0, bytemuck::cast_slice(&[lights])); + 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::(), 0u32); + + let offset = mem::size_of::() * 4 + mem::size_of::() * v as usize; + queue.write_buffer(&self.buffer, offset as _, bytemuck::cast_slice(zeros.as_slice())); + + self.dead_indices.push_back(v); + } + + // update the amount of lights, then the array of lights + queue.write_buffer(&self.buffer, 0, bytemuck::cast_slice(&[self.current_light_idx as u32])); } } @@ -227,8 +271,8 @@ pub(crate) struct LightUniform { pub intensity: f32, pub smoothness: f32, - pub spot_cutoff: f32, - pub spot_outer_cutoff: f32, + pub spot_cutoff_rad: f32, + pub spot_outer_cutoff_rad: f32, } impl LightUniform { @@ -244,8 +288,8 @@ impl LightUniform { intensity: light.intensity, smoothness: light.smoothness, - spot_cutoff: 0.0, - spot_outer_cutoff: 0.0, + spot_cutoff_rad: 0.0, + spot_outer_cutoff_rad: 0.0, } } @@ -262,190 +306,27 @@ impl LightUniform { intensity: light.intensity, smoothness: 0.0, - spot_cutoff: 0.0, - spot_outer_cutoff: 0.0, + spot_cutoff_rad: 0.0, + spot_outer_cutoff_rad: 0.0, } } // Create the SpotLightUniform from an ECS bundle - /* pub fn from_bundle(light: &SpotLight, transform: &Transform) -> Self { + pub fn from_spot_light_bundle(light: &SpotLight, transform: &Transform) -> Self { Self { + light_type: LightType::Spotlight as u32, + enabled: light.enabled as u32, position: transform.translation, - _padding: 0, direction: transform.forward(), - _padding2: 0, color: light.color, - cutoff: light.cutoff.to_radians().cos(), - outer_cutoff: light.outer_cutoff.to_radians().cos(), - constant: light.constant, - linear: light.linear, - quadratic: light.quadratic, - ambient: light.ambient, - diffuse: light.diffuse, - specular: light.specular, - _padding3: 0, - } - } */ -} -#[repr(C)] -#[derive(Default, Debug, Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)] -pub struct LightsUniform { - light_count: u32, - _padding: [u32; 3], - data: [LightUniform; 10], // TODO: make this a dynamic length -} - -/* #[repr(C)] -#[derive(Default, Debug, Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)] -pub struct LightsUniform { - light_count: u32, - _padding: [u32; 3], - spot_lights: [SpotLightUniform; MAX_LIGHT_COUNT], - spot_light_count: u32, - _padding2: [u32; 3], - directional_light: DirectionalLightUniform, -} - -#[repr(C)] -#[derive(Default, Debug, Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)] -pub struct PointLightUniform { - /// The position of the light - /// vec4 is used here for gpu padding, w is ignored in the shader - pub position: glam::Vec4, - - /// The color of the light - /// vec4 is used here for gpu padding, w is ignored in the shader - pub color: glam::Vec4, - - /// The intensity of the light - /// This works by just multiplying the result of the lighting - /// calculations by this scalar - pub intensity: f32, - - /// The constant used in the quadratic attenuation calculation. Its best to leave this at 1.0 - pub constant: f32, - - /// The linear factor used in the quadratic attenuation calculation. - pub linear: f32, - - /// The quadratic factor used in the quadratic attenuation calculation. - pub quadratic: f32, - - pub ambient: f32, - pub diffuse: f32, - pub specular: f32, - - pub _padding: u32, -} - -impl PointLightUniform { - /// Create the PointLightUniform from an ECS bundle - pub fn from_bundle(light: &PointLight, transform: &Transform) -> Self { - Self { - position: glam::Vec4::new(transform.translation.x, transform.translation.y, transform.translation.z, 0.0), - //_padding: 0, - color: glam::Vec4::new(light.color.x, light.color.y, light.color.z, 0.0), - //_padding2: 0, + range: light.range, intensity: light.intensity, - constant: light.constant, - linear: light.linear, - quadratic: light.quadratic, + smoothness: light.smoothness, - ambient: light.ambient, - diffuse: light.diffuse, - specular: light.specular, - - _padding: 0, + spot_cutoff_rad: light.cutoff.to_radians(), + spot_outer_cutoff_rad: light.outer_cutoff.to_radians(), } } } -#[repr(C)] -#[derive(Default, Debug, Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)] -pub struct DirectionalLightUniform { - /// The direction of the light - pub direction: glam::Vec3, - // gpu padding - pub _padding: u32, - /// The color of the light - pub color: glam::Vec3, - // no padding is needed here since ambient acts as the padding - // that would usually be needed for the vec3 - - /// The scalar of the ambient light created by this caster. - pub ambient: f32, - /// The scalar of the diffuse light created by this caster. - pub diffuse: f32, - /// The scalar of the specular reflections created by this caster. - pub specular: f32, - - pub _padding2: [u32; 2], -} - -impl DirectionalLightUniform { - /// Create the DirectionalLightUniform from an ECS bundle - pub fn from_bundle(light: &DirectionalLight, transform: &Transform) -> Self { - //transform.forward() - Self { - direction: transform.forward(), - _padding: 0, - color: light.color, - ambient: light.ambient, - diffuse: light.diffuse, - specular: light.specular, - _padding2: [0; 2], - } - } -} - -#[repr(C)] -#[derive(Default, Debug, Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)] -pub(crate) struct SpotLightUniform { - pub position: glam::Vec3, - pub _padding: u32, - pub direction: glam::Vec3, - pub _padding2: u32, - pub color: glam::Vec3, - // no padding is needed here since cutoff acts as the padding - // that would usually be needed for the vec3 - - pub cutoff: f32, - pub outer_cutoff: f32, - - /// The constant used in the quadratic attenuation calculation. Its best to leave this at 1.0 - pub constant: f32, - - /// The linear factor used in the quadratic attenuation calculation. - pub linear: f32, - - /// The quadratic factor used in the quadratic attenuation calculation. - pub quadratic: f32, - - pub ambient: f32, - pub diffuse: f32, - pub specular: f32, - pub _padding3: u32, -} - -impl SpotLightUniform { - /// Create the SpotLightUniform from an ECS bundle - pub fn from_bundle(light: &SpotLight, transform: &Transform) -> Self { - Self { - position: transform.translation, - _padding: 0, - direction: transform.forward(), - _padding2: 0, - color: light.color, - cutoff: light.cutoff.to_radians().cos(), - outer_cutoff: light.outer_cutoff.to_radians().cos(), - constant: light.constant, - linear: light.linear, - quadratic: light.quadratic, - ambient: light.ambient, - diffuse: light.diffuse, - specular: light.specular, - _padding3: 0, - } - } -} */ \ No newline at end of file diff --git a/lyra-game/src/render/light/spotlight.rs b/lyra-game/src/render/light/spotlight.rs index c0eeed8..fa89c5a 100644 --- a/lyra-game/src/render/light/spotlight.rs +++ b/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, - pub quadratic: f32, +impl Default for SpotLight { + fn default() -> Self { + Self { + enabled: true, + color: glam::Vec3::new(1.0, 1.0, 1.0), + range: 1.0, + intensity: 1.0, + smoothness: 0.75, - pub ambient: f32, - pub diffuse: f32, - pub specular: f32, + cutoff: math::Angle::Degrees(45.0), + outer_cutoff: math::Angle::Degrees(45.0), + } + } } \ No newline at end of file diff --git a/lyra-game/src/render/light_cull_compute.rs b/lyra-game/src/render/light_cull_compute.rs index 98b742e..66fd876 100644 --- a/lyra-game/src/render/light_cull_compute.rs +++ b/lyra-game/src/render/light_cull_compute.rs @@ -83,7 +83,6 @@ impl LightCullCompute { label: Some("BGL_LightIndicesGrid"), }); - // TODO: resize texture when screen is resized let size = wgpu::Extent3d { width: workgroup_size.x, height: workgroup_size.y, @@ -168,7 +167,8 @@ impl LightCullCompute { source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(shader_src)), }); - let workgroup_size = glam::UVec2::new((screen_size.width as f32 / 16.0).ceil() as u32, (screen_size.height as f32 / 16.0).ceil() as u32); + let workgroup_size = glam::UVec2::new((screen_size.width as f32 / 16.0).ceil() as u32, + (screen_size.height as f32 / 16.0).ceil() as u32); let light_grid = Self::create_grid(&device, workgroup_size); let depth_tex_pair = depth_texture.create_bind_group(&device); @@ -205,7 +205,11 @@ impl LightCullCompute { pub fn update_screen_size(&mut self, size: PhysicalSize) { self.screen_size_buffer.write_buffer(&self.queue, 0, &[UVec2::new(size.width, size.height)]); - self.workgroup_size = glam::UVec2::new((size.width as f32 / 16.0).ceil() as u32, (size.height as f32 / 16.0).ceil() as u32); + self.workgroup_size = glam::UVec2::new((size.width as f32 / 16.0).ceil() as u32, + (size.height as f32 / 16.0).ceil() as u32); + + // I hate that the entire bind group is recreated on a resize but its the only way :( + self.light_indices_grid = Self::create_grid(&self.device, self.workgroup_size); } pub fn compute(&mut self, camera_buffers: &BufferWrapper, lights_buffers: &LightUniformBuffers, depth_texture: &RenderTexture) { @@ -228,6 +232,7 @@ impl LightCullCompute { pass.dispatch_workgroups(self.workgroup_size.x, self.workgroup_size.y, 1); } + self.queue.submit(std::iter::once(encoder.finish())); self.device.poll(wgpu::Maintain::Wait); self.cleanup(); diff --git a/lyra-game/src/render/shaders/base.wgsl b/lyra-game/src/render/shaders/base.wgsl index aa907a7..7b97ddf 100755 --- a/lyra-game/src/render/shaders/base.wgsl +++ b/lyra-game/src/render/shaders/base.wgsl @@ -238,5 +238,50 @@ fn blinn_phong_point_light(world_pos: vec3, world_norm: vec3, point_li } fn blinn_phong_spot_light(world_pos: vec3, world_norm: vec3, spot_light: Light, material: Material, specular_factor: vec3) -> vec3 { - return vec3(0.0); // TODO + let light_color = spot_light.color; + let light_pos = spot_light.position; + let camera_view_pos = u_camera.position; + + let light_dir = normalize(spot_light.position - world_pos); + + var ambient_color = light_color * material.ambient.xyz * material.diffuse.xyz; + + //// diffuse //// + //let light_dir = normalize(light_pos - world_pos); + + let diffuse_strength = max(dot(world_norm, light_dir), 0.0); + var diffuse_color = light_color * (diffuse_strength * material.diffuse.xyz); + //// end of diffuse //// + + //// specular //// + let view_dir = normalize(camera_view_pos - world_pos); + let half_dir = normalize(view_dir + light_dir); + + let specular_strength = pow(max(dot(world_norm, half_dir), 0.0), material.shininess); + var specular_color = specular_strength * (light_color * specular_factor); + //// end of specular //// + + //// spot light soft edges //// + let theta = dot(light_dir, normalize(-spot_light.direction)); + let epsilon = spot_light.spot_cutoff - spot_light.spot_outer_cutoff; + let intensity = clamp((theta - spot_light.spot_outer_cutoff) / epsilon, 0.0, 1.0); + //diffuse_color *= intensity; + //specular_color *= intensity; + //// end of spot light soft edges //// + + //// spot light attenuation //// + let distance = length(light_pos - world_pos); + let attenuation = calc_attenuation(spot_light, distance); + + ambient_color *= attenuation * intensity; + diffuse_color *= attenuation * intensity; + specular_color *= attenuation * intensity; + //// end of spot light attenuation //// + + + return /*ambient_color +*/ diffuse_color + specular_color; +} + +fn calc_attenuation(light: Light, distance: f32) -> f32 { + return 1.0 - smoothstep(light.range * light.smoothness, light.range, distance); } \ No newline at end of file diff --git a/lyra-game/src/render/shaders/light_cull.comp.wgsl b/lyra-game/src/render/shaders/light_cull.comp.wgsl index d8e481b..dfe5a7a 100644 --- a/lyra-game/src/render/shaders/light_cull.comp.wgsl +++ b/lyra-game/src/render/shaders/light_cull.comp.wgsl @@ -33,6 +33,13 @@ struct Lights { data: array, }; +struct Cone { + tip: vec3, + height: f32, + direction: vec3, + radius: f32, +} + var wg_min_depth: atomic; var wg_max_depth: atomic; var wg_light_index_start: atomic; @@ -160,13 +167,22 @@ fn cs_main( && sphere_inside_frustrum(wg_frustum_planes, position, radius)) { // TODO: add the light to the transparent geometry list - add_light(light_index); - if (!sphere_inside_plane(position, radius, wg_frustum_planes[4])) { - + add_light(light_index); + } + } else if (light.light_ty == LIGHT_TY_SPOT) { + let dir_vs = u_camera.view * vec4(light.direction, 1.0); + let cone_radius = tan(light.spot_cutoff) * light.range; + let cone = Cone(position, radius, dir_vs.xyz, cone_radius); + + if (cone_inside_frustum(cone, wg_frustum_planes)) { + // TODO: add the light to the transparent geometry list + + if (!cone_inside_plane(cone, wg_frustum_planes[4])) { + add_light(light_index); + } } } - // TODO: spotlights } } @@ -258,4 +274,45 @@ fn compute_plane(p0: vec3, p1: vec3, p2: vec3) -> vec4 { plane.w = dot(plane.xyz, p0); return plane; +} + +fn point_inside_plane(point: vec3, plane: vec4) -> bool { + return dot(plane.xyz, point) - plane.w < 0.0; +} + +/// Check to see if a cone if fully behind (inside the negative halfspace of) a plane. +/// +/// Source: Real-time collision detection, Christer Ericson (2005) +/// (https://www.3dgep.com/forward-plus/#light-culling-compute-shader) +fn cone_inside_plane(cone: Cone, plane: vec4) -> bool { + // Compute the farthest point on the end of the cone to the positive space of the plane. + let m = cross(cross(plane.xyz, cone.direction), cone.direction); + let farthest = cone.tip + cone.direction * cone.height - m * cone.radius; + + // The cone is in the negative halfspace of the plane if the tip of the cone, + // and the farthest point on the end of the cone are inside the negative halfspace + // of the plane. + return point_inside_plane(cone.tip, plane) && point_inside_plane(farthest, plane); +} + +fn cone_inside_frustum(cone: Cone, frustum: array, 6>) -> bool { + //let near_plane = frustum[4]; + //let far_plane = frustum[5]; + + // check near and far clipping planes first + //if (cone_inside_plane(cone, near_plane) || cone_inside_plane(cone, far_plane)) { + // return false; + //} + + // to be able to index this array with a non-const value, + // it must be defined as a var + var frustum_v = frustum; + + for (var i = 0u; i < 4u; i++) { + if (cone_inside_plane(cone, frustum_v[i])) { + return false; + } + } + + return true; } \ No newline at end of file diff --git a/shell.nix b/shell.nix index d0c65e9..98db997 100755 --- a/shell.nix +++ b/shell.nix @@ -15,13 +15,19 @@ in mold udev lua5_4_compat - (nixpkgs.rustChannelOf { rustToolchain = ./rust-toolchain.toml; }).rust + ((nixpkgs.rustChannelOf { rustToolchain = ./rust-toolchain.toml; }).rust.override { + extensions = [ + "rust-src" + "rust-analysis" + ]; + }) ]; buildInputs = [ udev alsa-lib libGL gcc vulkan-loader vulkan-headers vulkan-tools xorg.libX11 xorg.libXcursor xorg.libXi xorg.libXrandr # To use the x11 feature libxkbcommon wayland # To use the wayland feature + ]; LD_LIBRARY_PATH = lib.makeLibraryPath buildInputs; } -- 2.40.1 From e2844a11a606ebcf3d59aecc8b68ef527392c863 Mon Sep 17 00:00:00 2001 From: SeanOMik Date: Wed, 20 Mar 2024 19:03:39 -0400 Subject: [PATCH 13/15] render: create toggleable debug light cull view --- examples/testbed/src/free_fly_camera.rs | 4 ++++ examples/testbed/src/main.rs | 11 ++++++++--- lyra-game/src/input/buttons.rs | 15 ++++++++++++++- lyra-game/src/input/system.rs | 5 ++++- lyra-game/src/render/camera.rs | 12 +++++++----- lyra-game/src/render/shaders/base.wgsl | 13 ++++++++++--- lyra-game/src/scene/camera.rs | 2 ++ 7 files changed, 49 insertions(+), 13 deletions(-) diff --git a/examples/testbed/src/free_fly_camera.rs b/examples/testbed/src/free_fly_camera.rs index b6f404b..fa05a8b 100644 --- a/examples/testbed/src/free_fly_camera.rs +++ b/examples/testbed/src/free_fly_camera.rs @@ -86,6 +86,10 @@ pub fn free_fly_camera_controller(delta_time: Res, handler: Res InputButtons { 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()); + }, + _ => {}, + } + } + } } \ No newline at end of file diff --git a/lyra-game/src/input/system.rs b/lyra-game/src/input/system.rs index fb04a0f..b41a939 100755 --- a/lyra-game/src/input/system.rs +++ b/lyra-game/src/input/system.rs @@ -102,12 +102,15 @@ impl crate::ecs::system::System for InputSystem { let queue = world.try_get_resource_mut::() .and_then(|q| q.read_events::()); + let mut e = world.get_resource_or_else(InputButtons::::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); } diff --git a/lyra-game/src/render/camera.rs b/lyra-game/src/render/camera.rs index 882f7c5..3b5ab46 100755 --- a/lyra-game/src/render/camera.rs +++ b/lyra-game/src/render/camera.rs @@ -46,7 +46,8 @@ pub struct CameraUniform { pub view_projection: glam::Mat4, /// The position of the camera pub position: glam::Vec3, - _padding: u32, + pub tile_debug: u32, + //_padding: [u8; 3], } @@ -57,7 +58,8 @@ impl Default for CameraUniform { inverse_projection: glam::Mat4::IDENTITY, view_projection: glam::Mat4::IDENTITY, position: Default::default(), - _padding: 0, + tile_debug: 0, + //_padding: 0, } } } @@ -69,7 +71,7 @@ impl CameraUniform { inverse_projection, view_projection, position, - _padding: 0 + tile_debug: 0 } } } @@ -128,7 +130,7 @@ impl RenderCamera { inverse_projection: proj.inverse(), view_projection: self.view_proj, position, - _padding: 0, + tile_debug: camera.debug as u32, } }, CameraProjectionMode::Orthographic => { @@ -150,7 +152,7 @@ impl RenderCamera { inverse_projection: proj.inverse(), view_projection: self.view_proj, position, - _padding: 0, + tile_debug: camera.debug as u32, } }, } diff --git a/lyra-game/src/render/shaders/base.wgsl b/lyra-game/src/render/shaders/base.wgsl index 7b97ddf..f91291e 100755 --- a/lyra-game/src/render/shaders/base.wgsl +++ b/lyra-game/src/render/shaders/base.wgsl @@ -24,6 +24,7 @@ struct CameraUniform { inverse_projection: mat4x4, view_projection: mat4x4, position: vec3, + tile_debug: u32, }; struct Light { @@ -127,7 +128,10 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4 { return vec4(light_object_res, object_color.a);*/ - + if (u_camera.tile_debug == 1u) { + return debug_grid(in); + } + let tile_index = vec2(floor(in.clip_position.xy / 16.0)); let tile: vec2 = textureLoad(t_light_grid, tile_index).xy; @@ -168,11 +172,14 @@ fn debug_grid(in: VertexOutput) -> vec4 { let ta: bool = x < 0.05 || y < 0.05; let tb: bool = x > 0.95 || y > 0.95; - if ( ta || tb ) { + let ratio = f32(tile.y) / f32(u_lights.light_count); + return vec4(ratio, ratio, ratio, 1.0); + + /* if ( ta || tb ) { return vec4(0.0, 0.0, 0.0, 1.0); } else { return vec4(f32(tile_index.x) / 50.0, f32(tile_index.y) / 38.0, 0.0, 1.0); - } + } */ } fn blinn_phong_dir_light(world_pos: vec3, world_norm: vec3, dir_light: Light, material: Material, specular_factor: vec3) -> vec3 { diff --git a/lyra-game/src/scene/camera.rs b/lyra-game/src/scene/camera.rs index 36137a5..0808699 100755 --- a/lyra-game/src/scene/camera.rs +++ b/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, } } } -- 2.40.1 From 0f11fe2e6d2f024868f5fd14b446924bface8873 Mon Sep 17 00:00:00 2001 From: SeanOMik Date: Fri, 22 Mar 2024 10:46:52 -0400 Subject: [PATCH 14/15] render: fix spot light culling --- examples/testbed/src/main.rs | 1 + lyra-game/src/render/camera.rs | 7 +- lyra-game/src/render/shaders/base.wgsl | 32 ++-- .../src/render/shaders/light_cull.comp.wgsl | 161 ++++++++++-------- 4 files changed, 114 insertions(+), 87 deletions(-) diff --git a/examples/testbed/src/main.rs b/examples/testbed/src/main.rs index 5b2334f..865ab3d 100644 --- a/examples/testbed/src/main.rs +++ b/examples/testbed/src/main.rs @@ -182,6 +182,7 @@ async fn main() { 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( diff --git a/lyra-game/src/render/camera.rs b/lyra-game/src/render/camera.rs index 3b5ab46..3c12dd4 100755 --- a/lyra-game/src/render/camera.rs +++ b/lyra-game/src/render/camera.rs @@ -44,6 +44,7 @@ pub struct CameraUniform { 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, @@ -57,6 +58,7 @@ impl Default for CameraUniform { view: glam::Mat4::IDENTITY, inverse_projection: glam::Mat4::IDENTITY, view_projection: glam::Mat4::IDENTITY, + projection: glam::Mat4::IDENTITY, position: Default::default(), tile_debug: 0, //_padding: 0, @@ -65,11 +67,12 @@ impl Default for CameraUniform { } impl CameraUniform { - pub fn new(view: glam::Mat4, inverse_projection: glam::Mat4, view_projection: glam::Mat4, position: glam::Vec3) -> Self { + 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 } @@ -129,6 +132,7 @@ impl RenderCamera { view, inverse_projection: proj.inverse(), view_projection: self.view_proj, + projection: proj, position, tile_debug: camera.debug as u32, } @@ -151,6 +155,7 @@ impl RenderCamera { view, inverse_projection: proj.inverse(), view_projection: self.view_proj, + projection: proj, position, tile_debug: camera.debug as u32, } diff --git a/lyra-game/src/render/shaders/base.wgsl b/lyra-game/src/render/shaders/base.wgsl index f91291e..d059aa0 100755 --- a/lyra-game/src/render/shaders/base.wgsl +++ b/lyra-game/src/render/shaders/base.wgsl @@ -6,6 +6,8 @@ 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, @location(1) tex_coords: vec2, @@ -23,6 +25,7 @@ struct CameraUniform { view: mat4x4, inverse_projection: mat4x4, view_projection: mat4x4, + projection: mat4x4, position: vec3, tile_debug: u32, }; @@ -132,13 +135,17 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4 { return debug_grid(in); } - let tile_index = vec2(floor(in.clip_position.xy / 16.0)); - let tile: vec2 = textureLoad(t_light_grid, tile_index).xy; - let object_color: vec4 = textureSample(t_diffuse, s_diffuse, in.tex_coords); let specular_color: vec3 = textureSample(t_specular, s_specular, in.tex_coords).xyz; var light_res = vec3(0.0); + if (object_color.a < ALPHA_CUTOFF) { + discard; + } + + let tile_index = vec2(floor(in.clip_position.xy / 16.0)); + let tile: vec2 = textureLoad(t_light_grid, tile_index).xy; + let light_offset = tile.x; let light_count = tile.y; @@ -269,20 +276,19 @@ fn blinn_phong_spot_light(world_pos: vec3, world_norm: vec3, spot_ligh //// end of specular //// //// spot light soft edges //// - let theta = dot(light_dir, normalize(-spot_light.direction)); - let epsilon = spot_light.spot_cutoff - spot_light.spot_outer_cutoff; - let intensity = clamp((theta - spot_light.spot_outer_cutoff) / epsilon, 0.0, 1.0); - //diffuse_color *= intensity; - //specular_color *= intensity; + let min_cos = cos(spot_light.spot_cutoff); + 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 attenuation //// let distance = length(light_pos - world_pos); let attenuation = calc_attenuation(spot_light, distance); - ambient_color *= attenuation * intensity; - diffuse_color *= attenuation * intensity; - specular_color *= attenuation * intensity; + ambient_color *= attenuation * spot_light.intensity * cone; + diffuse_color *= attenuation * spot_light.intensity * cone; + specular_color *= attenuation * spot_light.intensity * cone; //// end of spot light attenuation //// @@ -291,4 +297,8 @@ fn blinn_phong_spot_light(world_pos: vec3, world_norm: vec3, spot_ligh 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); } \ No newline at end of file diff --git a/lyra-game/src/render/shaders/light_cull.comp.wgsl b/lyra-game/src/render/shaders/light_cull.comp.wgsl index dfe5a7a..fd3552d 100644 --- a/lyra-game/src/render/shaders/light_cull.comp.wgsl +++ b/lyra-game/src/render/shaders/light_cull.comp.wgsl @@ -5,20 +5,25 @@ const LIGHT_TY_DIRECTIONAL = 0u; const LIGHT_TY_POINT = 1u; const LIGHT_TY_SPOT = 2u; +type vec2f = vec2; +type vec3f = vec3; +type vec4f = vec4; + struct CameraUniform { view: mat4x4, inverse_projection: mat4x4, - //projection: mat4x4, view_projection: mat4x4, - position: vec3, + projection: mat4x4, + position: vec3f, + tile_debug: u32, }; struct Light { - position: vec3, + position: vec3f, light_ty: u32, - direction: vec3, + direction: vec3f, enabled: u32, - color: vec3, + color: vec3f, range: f32, intensity: f32, @@ -34,16 +39,21 @@ struct Lights { }; struct Cone { - tip: vec3, + tip: vec3f, height: f32, - direction: vec3, + direction: vec3f, radius: f32, } +struct Plane { + normal: vec3f, + origin_distance: f32, +} + var wg_min_depth: atomic; var wg_max_depth: atomic; var wg_light_index_start: atomic; -var wg_frustum_planes: array, 6>; +var wg_frustum_planes: array; // index list of visible light sources for this tile var wg_visible_light_indices: array; @@ -107,41 +117,50 @@ fn cs_main( // Create the frustum planes that will be used for this time if (local_invocation_index == 0u) { - // Compute the 4 corner points on the far clipping plane to use as the frustum vertices. - var screen_space: array, 4>; + // 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); - // top left point - var temp: vec2 = workgroup_id.xy * BLOCK_SIZE; - screen_space[0] = vec4(f32(temp.x), f32(temp.y), -1.0, 1.0); + let tr_projection = transpose(u_camera.projection); - // top right point - var temp2 = vec2(f32(workgroup_id.x) + 1.0, f32(workgroup_id.y)) * f32(BLOCK_SIZE); - screen_space[1] = vec4(temp2.x, temp2.y, -1.0, 1.0); - - // bottom left point - temp2 = vec2(f32(workgroup_id.x), f32(workgroup_id.y) + 1.0) * f32(BLOCK_SIZE); - screen_space[2] = vec4(temp2.x, temp2.y, -1.0, 1.0); - - // bottom right point - temp2 = vec2(f32(workgroup_id.x) + 1.0, f32(workgroup_id.y) + 1.0) * f32(BLOCK_SIZE); - screen_space[3] = vec4(temp2.x, temp2.y, -1.0, 1.0); + var planes: array; - // convert screenspace to view space - var view_space: array, 4>; - for (var i = 0u; i < 4u; i++) { - view_space[i] = screen_to_view(screen_space[i]).xyz; + // 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); } - // View space eye is always at the origin - let eye_pos = vec3(0.0, 0.0, 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); + } - wg_frustum_planes[0] = compute_plane(eye_pos, view_space[2], view_space[0]); // left plane - wg_frustum_planes[1] = compute_plane(eye_pos, view_space[1], view_space[3]); // right plane - wg_frustum_planes[2] = compute_plane(eye_pos, view_space[0], view_space[1]); // top plane - wg_frustum_planes[3] = compute_plane(eye_pos, view_space[3], view_space[2]); // bottom plane + // 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); + } - wg_frustum_planes[4] = vec4(0.0, 0.0, -1.0, -min_depth); - wg_frustum_planes[5] = vec4(0.0, 0.0, 1.0, -max_depth); + // 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(); @@ -157,29 +176,27 @@ fn cs_main( let light = u_lights.data[light_index]; if (light.enabled == 1u) { - let position_vec4 = u_camera.view * vec4(light.position, 1.0); - let position = position_vec4.xyz; - let radius = light.range; + 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, radius)) { + && sphere_inside_frustrum(wg_frustum_planes, position_vs, light.range)) { // TODO: add the light to the transparent geometry list - if (!sphere_inside_plane(position, radius, wg_frustum_planes[4])) { + 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 * vec4(light.direction, 1.0); + 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, radius, dir_vs.xyz, cone_radius); + 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])) { - add_light(light_index); } } } @@ -220,7 +237,7 @@ fn add_light(light_index: u32) -> bool { return false; } -fn sphere_inside_frustrum(frustum: array, 6>, sphere_origin: vec3, radius: f32) -> bool { +fn sphere_inside_frustrum(frustum: array, 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; @@ -239,11 +256,11 @@ fn sphere_inside_frustrum(frustum: array, 6>, sphere_origin: vec3 /// /// 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, radius: f32, plane: vec4) -> bool { - return dot(plane.xyz, sphere_origin) - plane.w < -radius; +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: vec4) -> vec4 { +fn clip_to_view(clip: vec4f) -> vec4f { // view space position var view = u_camera.inverse_projection * clip; @@ -251,12 +268,12 @@ fn clip_to_view(clip: vec4) -> vec4 { return view / view.w; } -fn screen_to_view(screen: vec4) -> vec4 { +fn screen_to_view(screen: vec4f) -> vec4f { // convert to normalized texture coordinates let tex_coord = screen.xy / vec2(u_screen_size); // convert to clip space - let clip = vec4( vec2(tex_coord.x, 1.0 - tex_coord.y) * 2.0 - 1.0, screen.z, screen.w); + let clip = vec4f( vec2(tex_coord.x, 1.0 - tex_coord.y) * 2.0 - 1.0, screen.z, screen.w); return clip_to_view(clip); } @@ -264,52 +281,46 @@ fn screen_to_view(screen: vec4) -> vec4 { /// 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: vec3, p1: vec3, p2: vec3) -> vec4 { +fn compute_plane(p0: vec3f, p1: vec3f, p2: vec3f) -> Plane { let v0 = p1 - p0; let v2 = p2 - p0; - var plane = vec4(normalize(cross(v0, v2)), 0.0); + let normal = vec4f(normalize(cross(v0, v2)), 0.0); // find the distance to the origin - plane.w = dot(plane.xyz, p0); + let distance = dot(normal.xyz, p0); - return plane; + return Plane(normal.xyz, distance); } -fn point_inside_plane(point: vec3, plane: vec4) -> bool { - return dot(plane.xyz, point) - plane.w < 0.0; +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: vec4) -> bool { - // Compute the farthest point on the end of the cone to the positive space of the plane. - let m = cross(cross(plane.xyz, cone.direction), cone.direction); - let farthest = cone.tip + cone.direction * cone.height - m * cone.radius; +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(farthest, plane); + return point_inside_plane(cone.tip, plane) && point_inside_plane(furthest, plane); } -fn cone_inside_frustum(cone: Cone, frustum: array, 6>) -> bool { - //let near_plane = frustum[4]; - //let far_plane = frustum[5]; - - // check near and far clipping planes first - //if (cone_inside_plane(cone, near_plane) || cone_inside_plane(cone, far_plane)) { - // return false; - //} - - // to be able to index this array with a non-const value, - // it must be defined as a var - var frustum_v = frustum; - +fn cone_inside_frustum(cone: Cone, frustum: array) -> bool { + var frustum = frustum; for (var i = 0u; i < 4u; i++) { - if (cone_inside_plane(cone, frustum_v[i])) { + // TODO: better cone checking + if (sphere_inside_plane(cone.tip, cone.radius, frustum[i])) { return false; } } -- 2.40.1 From 763d51ae364b9adbe129da90c84ff7d80ffc1367 Mon Sep 17 00:00:00 2001 From: SeanOMik Date: Fri, 22 Mar 2024 22:55:22 -0400 Subject: [PATCH 15/15] move some stuff out of testbed into lyra-game --- examples/testbed/Cargo.toml | 3 +- examples/testbed/src/free_fly_camera.rs | 106 ------------------------ examples/testbed/src/main.rs | 85 ++++--------------- lyra-game/src/scene/free_fly_camera.rs | 99 +++++++++++++++++++--- 4 files changed, 103 insertions(+), 190 deletions(-) delete mode 100644 examples/testbed/src/free_fly_camera.rs diff --git a/examples/testbed/Cargo.toml b/examples/testbed/Cargo.toml index 79687cc..5a52c01 100644 --- a/examples/testbed/Cargo.toml +++ b/examples/testbed/Cargo.toml @@ -12,5 +12,4 @@ lyra-engine = { path = "../../", version = "0.0.1" } #lyra-ecs = { path = "../../lyra-ecs"} anyhow = "1.0.75" async-std = "1.12.0" -tracing = "0.1.37" -fps_counter = "2.0.0" \ No newline at end of file +tracing = "0.1.37" \ No newline at end of file diff --git a/examples/testbed/src/free_fly_camera.rs b/examples/testbed/src/free_fly_camera.rs deleted file mode 100644 index fa05a8b..0000000 --- a/examples/testbed/src/free_fly_camera.rs +++ /dev/null @@ -1,106 +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.5, - mouse_sensitivity: 0.9, - 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, handler: Res, 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); - } - - if handler.was_action_just_pressed("Debug") { - cam.debug = !cam.debug; - } - } - - 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, &[]); - } -} diff --git a/examples/testbed/src/main.rs b/examples/testbed/src/main.rs index 865ab3d..be558ba 100644 --- a/examples/testbed/src/main.rs +++ b/examples/testbed/src/main.rs @@ -1,29 +1,8 @@ -use std::{cell::Ref, ptr::NonNull}; +use std::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, 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"; -const ACTLBL_DEBUG: &str = "Debug"; - struct FixedTimestep { max_tps: u32, fixed_time: f32, @@ -254,52 +233,18 @@ async fn main() { Ok(()) }; - #[allow(unused_variables)] - let fps_system = |world: &mut World| -> anyhow::Result<()> { - let mut counter = world.get_resource_mut::(); + let camera_debug_plugin = move |game: &mut Game| { + let sys = |handler: Res, view: View<&mut CameraComponent>| -> anyhow::Result<()> { + 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}"); - - 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::(); - - 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()); */ - - /* let mut camera = CameraComponent::new_3d(); - camera.transform.translation += math::Vec3::new(0.0, 0.0, 5.5); - world.spawn(( camera, FreeFlyCamera::default() )); */ - //sys.with_system(fps_system); - - //game.with_system("fixed", sys, &[]); - //fps_plugin(game); + game.with_system("camera_debug_trigger", sys, &[]); }; let action_handler_plugin = |game: &mut Game| { @@ -312,7 +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(ACTLBL_DEBUG, Action::new(ActionKind::Button)) + .add_action("Debug", Action::new(ActionKind::Button)) .add_mapping(ActionMapping::builder(LayoutId::from(0), ActionMappingId::from(0)) .bind(ACTLBL_MOVE_FORWARD_BACKWARD, &[ @@ -343,7 +288,7 @@ async fn main() { ActionSource::Keyboard(KeyCode::E).into_binding_modifier(-1.0), ActionSource::Keyboard(KeyCode::Q).into_binding_modifier(1.0), ]) - .bind(ACTLBL_DEBUG, &[ + .bind("Debug", &[ ActionSource::Keyboard(KeyCode::B).into_binding(), ]) .finish() @@ -375,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; } diff --git a/lyra-game/src/scene/free_fly_camera.rs b/lyra-game/src/scene/free_fly_camera.rs index abfc6f4..b02e3f7 100644 --- a/lyra-game/src/scene/free_fly_camera.rs +++ b/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(), - fov: Angle::Degrees(45.0), - mode: CameraProjectionMode::Perspective, - speed: 1.5, + speed, + slow_speed_factor, + look_speed, + mouse_sensitivity, + look_with_keys, } } -} \ No newline at end of file +} + +pub fn free_fly_camera_controller(delta_time: Res, handler: Res, 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, &[]); + } +} -- 2.40.1