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: