Compare commits
No commits in common. "5c1ce809ff0f24502e97961323161484b0b6ac04" and "1818a0b48b33a14170cdc867f64c82d8df52c265" have entirely different histories.
5c1ce809ff
...
1818a0b48b
|
|
@ -25,8 +25,8 @@ impl Default for FreeFlyCamera {
|
|||
Self {
|
||||
speed: 4.0,
|
||||
slow_speed_factor: 0.25,
|
||||
look_speed: 0.5,
|
||||
mouse_sensitivity: 0.9,
|
||||
look_speed: 0.3,
|
||||
mouse_sensitivity: 1.0,
|
||||
look_with_keys: false,
|
||||
}
|
||||
}
|
||||
|
|
|
|||
|
|
@ -1,6 +1,6 @@
|
|||
use std::{cell::Ref, ptr::NonNull};
|
||||
|
||||
use lyra_engine::{assets::gltf::Gltf, change_tracker::Ct, ecs::{query::{QueryBorrow, ViewState}, system::{BatchedSystem, Criteria, CriteriaSchedule, IntoSystem}, Component, World}, game::Game, input::{Action, ActionHandler, ActionKind, ActionMapping, ActionMappingId, ActionSource, InputActionPlugin, KeyCode, LayoutId, MouseAxis, MouseInput}, math::{self, Quat, Transform, Vec3}, render::{light::{directional::DirectionalLight, PointLight, SpotLight}, window::{CursorGrabMode, WindowOptions}}, scene::CameraComponent, DeltaTime};
|
||||
use lyra_engine::{assets::gltf::Gltf, ecs::{query::{QueryBorrow, ViewState}, system::{BatchedSystem, Criteria, CriteriaSchedule, IntoSystem}, Component, World}, game::Game, input::{Action, ActionHandler, ActionKind, ActionMapping, ActionMappingId, ActionSource, InputActionPlugin, KeyCode, LayoutId, MouseAxis, MouseInput}, math::{self, Quat, Transform, Vec3}, render::light::{directional::DirectionalLight, SpotLight}, scene::CameraComponent, DeltaTime};
|
||||
use lyra_engine::assets::ResourceManager;
|
||||
|
||||
mod free_fly_camera;
|
||||
|
|
@ -93,16 +93,15 @@ async fn main() {
|
|||
//let diffuse_texture = resman.request::<Texture>("assets/happy-tree.png").unwrap();
|
||||
//let antique_camera_model = resman.request::<Model>("assets/AntiqueCamera.glb").unwrap();
|
||||
//let cube_model = resman.request::<Model>("assets/cube-texture-bin.glb").unwrap();
|
||||
let cube_gltf = resman.request::<Gltf>("assets/texture-sep/texture-sep.gltf").unwrap();
|
||||
/*let crate_gltf = resman.request::<Gltf>("assets/crate/crate.gltf").unwrap();
|
||||
/* let cube_gltf = resman.request::<Gltf>("assets/texture-sep/texture-sep.gltf").unwrap();
|
||||
let crate_gltf = resman.request::<Gltf>("assets/crate/crate.gltf").unwrap();
|
||||
|
||||
let separate_gltf = resman.request::<Gltf>("assets/pos-testing/child-node-cubes.glb").unwrap(); */
|
||||
//drop(resman);
|
||||
|
||||
cube_gltf.wait_recurse_dependencies_load();
|
||||
let cube_mesh = &cube_gltf.data_ref()
|
||||
|
||||
/* let cube_mesh = &cube_gltf.data_ref()
|
||||
.unwrap().meshes[0];
|
||||
/* let crate_mesh = &crate_gltf.data_ref()
|
||||
let crate_mesh = &crate_gltf.data_ref()
|
||||
.unwrap().meshes[0];
|
||||
|
||||
let separate_scene = &separate_gltf.data_ref()
|
||||
|
|
@ -152,33 +151,29 @@ async fn main() {
|
|||
));
|
||||
}
|
||||
|
||||
{
|
||||
//let mut light_tran = Transform::from_xyz(-3.5, 0.2, -4.5);
|
||||
//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((
|
||||
PointLight {
|
||||
color: Vec3::new(0.0, 0.0, 1.0),
|
||||
|
||||
intensity: 3.3,
|
||||
SpotLight {
|
||||
color: Vec3::new(1.0, 0.2, 0.2),
|
||||
cutoff: math::Angle::Degrees(12.5),
|
||||
outer_cutoff: math::Angle::Degrees(17.5),
|
||||
|
||||
constant: 1.0,
|
||||
linear: 0.09,
|
||||
quadratic: 0.032,
|
||||
|
||||
ambient: 0.2,
|
||||
diffuse: 1.0,
|
||||
specular: 1.3,
|
||||
linear: 0.007,
|
||||
quadratic: 0.0002,
|
||||
|
||||
ambient: 0.0,
|
||||
diffuse: 7.0,
|
||||
specular: 1.0,
|
||||
},
|
||||
Transform::new(
|
||||
Vec3::new(-5.0, 1.0, -1.28),
|
||||
Quat::IDENTITY,
|
||||
Vec3::new(0.25, 0.25, 0.25),
|
||||
),
|
||||
Transform::from(light_tran),
|
||||
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((
|
||||
|
|
|
|||
|
|
@ -5,7 +5,7 @@
|
|||
"commandLine": "",
|
||||
"environment": [
|
||||
],
|
||||
"executable": "/media/data_drive/Development/Rust/lyra-engine/target/debug/testbed",
|
||||
"executable": "/media/data_drive/Development/Rust/lyra-test/engine/target/debug/testbed",
|
||||
"inject": false,
|
||||
"numQueuedFrames": 1,
|
||||
"options": {
|
||||
|
|
@ -23,6 +23,6 @@
|
|||
"verifyBufferAccess": false
|
||||
},
|
||||
"queuedFrameCap": 5,
|
||||
"workingDir": "/media/data_drive/Development/Rust/lyra-engine/examples/testbed"
|
||||
"workingDir": "/media/data_drive/Development/Rust/lyra-test/engine/examples/testbed"
|
||||
}
|
||||
}
|
||||
|
|
|
|||
|
|
@ -350,7 +350,6 @@ impl Game {
|
|||
// done by prefix, so it includes all lyra subpackages
|
||||
.with_target("lyra", Level::DEBUG)
|
||||
.with_target("wgpu", Level::WARN)
|
||||
.with_target("winit", Level::DEBUG)
|
||||
.with_default(Level::INFO))
|
||||
.init();
|
||||
|
||||
|
|
|
|||
|
|
@ -6,14 +6,14 @@ use lyra_ecs::{Entity, Tick, World, query::{Entities, TickOf}};
|
|||
pub use point::*;
|
||||
pub use spotlight::*;
|
||||
|
||||
use std::{collections::{HashMap, VecDeque}, marker::PhantomData};
|
||||
use std::{collections::{VecDeque, HashMap}, marker::PhantomData};
|
||||
|
||||
use std::mem;
|
||||
|
||||
use crate::math::Transform;
|
||||
|
||||
use self::directional::DirectionalLight;
|
||||
|
||||
use super::render_buffer::BindGroupPair;
|
||||
|
||||
const MAX_LIGHT_COUNT: usize = 16;
|
||||
|
||||
/// A struct that stores a list of lights in a wgpu::Buffer.
|
||||
|
|
@ -101,7 +101,8 @@ impl<U: Default + bytemuck::Pod + bytemuck::Zeroable> LightBuffer<U> {
|
|||
|
||||
pub(crate) struct LightUniformBuffers {
|
||||
pub buffer: wgpu::Buffer,
|
||||
pub bind_group_pair: BindGroupPair,
|
||||
pub bindgroup_layout: wgpu::BindGroupLayout,
|
||||
pub bindgroup: wgpu::BindGroup,
|
||||
pub light_indexes: HashMap<Entity, u32>,
|
||||
pub current_light_idx: u32,
|
||||
}
|
||||
|
|
@ -158,7 +159,8 @@ impl LightUniformBuffers {
|
|||
|
||||
Self {
|
||||
buffer,
|
||||
bind_group_pair: BindGroupPair::new(bindgroup, bindgroup_layout),
|
||||
bindgroup_layout,
|
||||
bindgroup,
|
||||
light_indexes: Default::default(),
|
||||
current_light_idx: 0,
|
||||
}
|
||||
|
|
|
|||
|
|
@ -1,36 +1,47 @@
|
|||
use std::{borrow::Cow, mem, rc::Rc};
|
||||
use std::{borrow::Cow, mem, num::NonZeroU32, ptr::NonNull, rc::Rc};
|
||||
|
||||
use glam::UVec2;
|
||||
use wgpu::{util::DeviceExt, ComputePipeline};
|
||||
use winit::dpi::PhysicalSize;
|
||||
|
||||
use super::{light::LightUniformBuffers, render_buffer::{BindGroupPair, BufferWrapper}, texture::RenderTexture};
|
||||
use super::{light::LightUniformBuffers, render_buffer::BufferWrapper, renderer::RenderPass, texture::RenderTexture};
|
||||
|
||||
pub(crate) struct LightIndicesGridBuffer {
|
||||
struct LightIndicesGridBuffer {
|
||||
indices_buffer: wgpu::Buffer,
|
||||
grid_texture: wgpu::Texture,
|
||||
grid_texture_view: wgpu::TextureView,
|
||||
pub bg_pair: BindGroupPair,
|
||||
bind_group_layout: wgpu::BindGroupLayout,
|
||||
bind_group: wgpu::BindGroup,
|
||||
}
|
||||
|
||||
pub(crate) struct LightCullCompute {
|
||||
device: Rc<wgpu::Device>,
|
||||
queue: Rc<wgpu::Queue>,
|
||||
pipeline: ComputePipeline,
|
||||
pub light_indices_grid: LightIndicesGridBuffer,
|
||||
lights: NonNull<LightUniformBuffers>,
|
||||
camera: NonNull<BufferWrapper>,
|
||||
light_indices_grid: LightIndicesGridBuffer,
|
||||
screen_size_buffer: BufferWrapper,
|
||||
workgroup_size: glam::UVec2,
|
||||
depth_tex: NonNull<RenderTexture>,
|
||||
}
|
||||
|
||||
impl LightCullCompute {
|
||||
fn create_grid(device: &wgpu::Device, screen_size: PhysicalSize<u32>, workgroup_size: glam::UVec2) -> LightIndicesGridBuffer {
|
||||
let mut contents = Vec::<u8>::new();
|
||||
let contents_len = workgroup_size.x * workgroup_size.y * mem::size_of::<u8>() as u32;
|
||||
contents.resize(contents_len as _, 0);
|
||||
fn create_grid(device: &wgpu::Device, screen_size: PhysicalSize<u32>) -> LightIndicesGridBuffer {
|
||||
let limits = device.limits();
|
||||
//let max_buffer_sizes = (limits.max_uniform_buffer_binding_size as u64) / 2;
|
||||
|
||||
/* let light_indices_buffer = device.create_buffer(
|
||||
&wgpu::BufferDescriptor {
|
||||
label: Some("B_LightIndices"),
|
||||
usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
|
||||
size: (mem::size_of::<u32>() * 16 * 16) as u64,
|
||||
mapped_at_creation: false,
|
||||
}
|
||||
); */
|
||||
|
||||
let light_indices_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
|
||||
label: Some("B_LightIndices"),
|
||||
contents: &contents,
|
||||
contents: &[0; mem::size_of::<u32>() * 16 * 16],
|
||||
usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
|
||||
});
|
||||
|
||||
|
|
@ -117,11 +128,12 @@ impl LightCullCompute {
|
|||
indices_buffer: light_indices_buffer,
|
||||
grid_texture,
|
||||
grid_texture_view,
|
||||
bg_pair: BindGroupPair::new(light_indices_bg, light_indices_bg_layout),
|
||||
bind_group_layout: light_indices_bg_layout,
|
||||
bind_group: light_indices_bg,
|
||||
}
|
||||
}
|
||||
|
||||
pub fn new(device: Rc<wgpu::Device>, queue: Rc<wgpu::Queue>, screen_size: PhysicalSize<u32>, lights_buffers: &LightUniformBuffers, camera_buffers: &BufferWrapper, depth_texture: &mut RenderTexture) -> Self {
|
||||
pub fn new(device: Rc<wgpu::Device>, queue: Rc<wgpu::Queue>, screen_size: PhysicalSize<u32>, lights: &LightUniformBuffers, camera_buffers: &BufferWrapper, depth_texture: &mut RenderTexture) -> Self {
|
||||
let screen_size_buffer = BufferWrapper::builder()
|
||||
.buffer_usage(wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST)
|
||||
.label_prefix("ScreenSize")
|
||||
|
|
@ -136,18 +148,18 @@ impl LightCullCompute {
|
|||
source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(shader_src)),
|
||||
});
|
||||
|
||||
let workgroup_size = glam::UVec2::new((screen_size.width as f32 / 16.0).ceil() as u32, (screen_size.height as f32 / 16.0).ceil() as u32);
|
||||
let light_grid = Self::create_grid(&device, screen_size, workgroup_size);
|
||||
|
||||
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_buffers.bind_group_pair.layout,
|
||||
&light_grid.bg_pair.layout,
|
||||
camera_buffers.bindgroup_layout().unwrap(),
|
||||
&lights.bindgroup_layout,
|
||||
&light_grid.bind_group_layout,
|
||||
screen_size_buffer.bindgroup_layout().unwrap(),
|
||||
],
|
||||
push_constant_ranges: &[],
|
||||
|
|
@ -164,19 +176,22 @@ impl LightCullCompute {
|
|||
device,
|
||||
queue,
|
||||
pipeline,
|
||||
lights: NonNull::from(lights),
|
||||
camera: NonNull::from(camera_buffers),
|
||||
light_indices_grid: light_grid,
|
||||
screen_size_buffer,
|
||||
workgroup_size,
|
||||
depth_tex: NonNull::from(depth_texture),
|
||||
}
|
||||
}
|
||||
|
||||
pub fn update_screen_size(&mut self, size: PhysicalSize<u32>) {
|
||||
pub fn update_screen_size(&self, size: PhysicalSize<u32>) {
|
||||
self.screen_size_buffer.write_buffer(&self.queue, 0,
|
||||
&[UVec2::new(size.width, size.height)]);
|
||||
self.workgroup_size = glam::UVec2::new((size.width as f32 / 16.0).ceil() as u32, (size.height as f32 / 16.0).ceil() as u32);
|
||||
}
|
||||
|
||||
pub fn compute(&mut self, camera_buffers: &BufferWrapper, lights_buffers: &LightUniformBuffers, depth_texture: &RenderTexture) {
|
||||
pub fn compute(&mut self) {
|
||||
//self.queue.write_buffer(&self.light_indices_grid.indices_buffer, 0, &[0; mem::size_of::<u32>() * 16 * 16]);
|
||||
|
||||
let mut encoder = self.device.create_command_encoder(&wgpu::CommandEncoderDescriptor {
|
||||
label: Some("LightCullCompute"),
|
||||
});
|
||||
|
|
@ -187,14 +202,20 @@ impl LightCullCompute {
|
|||
});
|
||||
|
||||
pass.set_pipeline(&self.pipeline);
|
||||
|
||||
let depth = unsafe { self.depth_tex.as_ref() };
|
||||
pass.set_bind_group(0, depth.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, &[]);
|
||||
let cam = unsafe { self.camera.as_ref() };
|
||||
pass.set_bind_group(1, cam.bindgroup(), &[]);
|
||||
|
||||
let lights = unsafe { self.lights.as_ref() };
|
||||
pass.set_bind_group(2, &lights.bindgroup, &[]);
|
||||
|
||||
pass.set_bind_group(3, &self.light_indices_grid.bind_group, &[]);
|
||||
pass.set_bind_group(4, self.screen_size_buffer.bindgroup(), &[]);
|
||||
|
||||
pass.dispatch_workgroups(self.workgroup_size.x, self.workgroup_size.y, 1);
|
||||
pass.dispatch_workgroups(16, 16, 1);
|
||||
}
|
||||
self.queue.submit(std::iter::once(encoder.finish()));
|
||||
self.device.poll(wgpu::Maintain::Wait);
|
||||
|
|
|
|||
|
|
@ -1,4 +1,4 @@
|
|||
use std::rc::Rc;
|
||||
use std::sync::Arc;
|
||||
|
||||
use lyra_resource::{ResHandle, Texture};
|
||||
|
||||
|
|
@ -11,7 +11,7 @@ pub struct MaterialSpecular {
|
|||
pub color_texture: Option<RenderTexture>,
|
||||
}
|
||||
|
||||
fn texture_to_render(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: &Rc<wgpu::BindGroupLayout>, i: &Option<ResHandle<Texture>>) -> Option<RenderTexture> {
|
||||
fn texture_to_render(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: &Arc<wgpu::BindGroupLayout>, i: &Option<ResHandle<Texture>>) -> Option<RenderTexture> {
|
||||
if let Some(tex) = i {
|
||||
RenderTexture::from_resource(device, queue, bg_layout.clone(), tex, None).ok()
|
||||
} else {
|
||||
|
|
@ -20,7 +20,7 @@ fn texture_to_render(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: &Rc<
|
|||
}
|
||||
|
||||
impl MaterialSpecular {
|
||||
pub fn from_resource(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Rc<wgpu::BindGroupLayout>, value: &lyra_resource::gltf::Specular) -> Self {
|
||||
pub fn from_resource(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Arc<wgpu::BindGroupLayout>, value: &lyra_resource::gltf::Specular) -> Self {
|
||||
let tex = texture_to_render(device, queue, &bg_layout, &value.texture);
|
||||
let color_tex = texture_to_render(device, queue, &bg_layout, &value.color_texture);
|
||||
|
||||
|
|
@ -45,7 +45,7 @@ pub struct Material {
|
|||
}
|
||||
|
||||
impl Material {
|
||||
pub fn from_resource(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Rc<wgpu::BindGroupLayout>, value: &lyra_resource::gltf::Material) -> Self {
|
||||
pub fn from_resource(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Arc<wgpu::BindGroupLayout>, value: &lyra_resource::gltf::Material) -> Self {
|
||||
let diffuse_texture = texture_to_render(device, queue, &bg_layout, &value.base_color_texture);
|
||||
|
||||
let specular = value.specular.as_ref().map(|s| MaterialSpecular::from_resource(device, queue, bg_layout.clone(), s));
|
||||
|
|
|
|||
|
|
@ -1,4 +1,4 @@
|
|||
use std::{num::NonZeroU32, rc::Rc};
|
||||
use std::{num::NonZeroU32, ops::Deref, sync::Arc};
|
||||
|
||||
use wgpu::util::DeviceExt;
|
||||
|
||||
|
|
@ -23,13 +23,13 @@ impl RenderBuffer {
|
|||
|
||||
pub struct BindGroupPair {
|
||||
pub bindgroup: wgpu::BindGroup,
|
||||
pub layout: Rc<wgpu::BindGroupLayout>,
|
||||
pub layout: Arc<wgpu::BindGroupLayout>,
|
||||
}
|
||||
|
||||
impl BindGroupPair {
|
||||
pub fn create_bind_group(device: &wgpu::Device, layout: Rc<wgpu::BindGroupLayout>, entries: &[wgpu::BindGroupEntry<'_>]) -> Self {
|
||||
pub fn new_from_layout(device: &wgpu::Device, layout: Arc<wgpu::BindGroupLayout>, entries: &[wgpu::BindGroupEntry<'_>]) -> Self {
|
||||
let bindgroup = device.create_bind_group(&wgpu::BindGroupDescriptor {
|
||||
layout: &layout,
|
||||
layout: layout.as_ref(),
|
||||
entries,
|
||||
label: None,
|
||||
});
|
||||
|
|
@ -43,7 +43,7 @@ impl BindGroupPair {
|
|||
pub fn new(bindgroup: wgpu::BindGroup, layout: wgpu::BindGroupLayout) -> Self {
|
||||
Self {
|
||||
bindgroup,
|
||||
layout: Rc::new(layout),
|
||||
layout: Arc::new(layout),
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
@ -92,7 +92,7 @@ impl BufferWrapper {
|
|||
///
|
||||
/// Returns None if this buffer object was not provided a bindgroup.
|
||||
pub fn bindgroup_layout(&self) -> Option<&wgpu::BindGroupLayout> {
|
||||
self.bindgroup_pair.as_ref().map(|bg| &*bg.layout)
|
||||
self.bindgroup_pair.as_ref().map(|bg| bg.layout.deref())
|
||||
}
|
||||
|
||||
/// Queue's the data to be written to `buffer` starting at `offset`.
|
||||
|
|
@ -272,7 +272,7 @@ impl BufferWrapperBuilder {
|
|||
],
|
||||
label: self.format_label("BGL_").as_deref(),
|
||||
});
|
||||
//let bg_layout = Arc::new(bg_layout);
|
||||
let bg_layout = Arc::new(bg_layout);
|
||||
|
||||
|
||||
let bg = device.create_bind_group(&wgpu::BindGroupDescriptor {
|
||||
|
|
@ -288,7 +288,7 @@ impl BufferWrapperBuilder {
|
|||
|
||||
BindGroupPair {
|
||||
bindgroup: bg,
|
||||
layout: Rc::new(bg_layout),
|
||||
layout: bg_layout,
|
||||
}
|
||||
}
|
||||
};
|
||||
|
|
|
|||
|
|
@ -59,7 +59,7 @@ struct MeshBufferStorage {
|
|||
|
||||
//#[allow(dead_code)]
|
||||
//render_texture: Option<RenderTexture>,
|
||||
material: Option<Rc<Material>>,
|
||||
material: Option<Material>,
|
||||
|
||||
// The index of the transform for this entity.
|
||||
// The tuple is structured like this: (transform index, index of transform inside the buffer)
|
||||
|
|
@ -88,7 +88,6 @@ pub struct BasicRenderer {
|
|||
pub render_jobs: VecDeque<RenderJob>,
|
||||
|
||||
mesh_buffers: HashMap<uuid::Uuid, MeshBufferStorage>, // TODO: clean up left over buffers from deleted entities/components
|
||||
material_buffers: HashMap<uuid::Uuid, Rc<Material>>,
|
||||
entity_meshes: HashMap<Entity, uuid::Uuid>,
|
||||
entity_last_transforms: HashMap<Entity, CachedTransform>,
|
||||
|
||||
|
|
@ -100,7 +99,7 @@ pub struct BasicRenderer {
|
|||
camera_buffer: BufferWrapper,
|
||||
//camera_bind_group: wgpu::BindGroup,
|
||||
|
||||
bgl_texture: Rc<BindGroupLayout>,
|
||||
bgl_texture: Arc<BindGroupLayout>,
|
||||
default_texture: RenderTexture,
|
||||
depth_buffer_texture: RenderTexture,
|
||||
|
||||
|
|
@ -173,7 +172,7 @@ impl BasicRenderer {
|
|||
};
|
||||
surface.configure(&device, &config);
|
||||
|
||||
let bgl_texture = Rc::new(RenderTexture::create_layout(&device));
|
||||
let bgl_texture = Arc::new(RenderTexture::create_layout(&device));
|
||||
|
||||
let shader_src = include_str!("shaders/base.wgsl");
|
||||
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
|
||||
|
|
@ -224,7 +223,6 @@ impl BasicRenderer {
|
|||
render_pipelines: HashMap::new(),
|
||||
render_jobs: VecDeque::new(),
|
||||
mesh_buffers: HashMap::new(),
|
||||
material_buffers: HashMap::new(),
|
||||
entity_meshes: HashMap::new(),
|
||||
|
||||
render_limits,
|
||||
|
|
@ -249,10 +247,8 @@ impl BasicRenderer {
|
|||
vec![super::vertex::Vertex::desc(),],
|
||||
vec![&s.bgl_texture, &s.transform_buffers.bindgroup_layout,
|
||||
s.camera_buffer.bindgroup_layout().unwrap(),
|
||||
&s.light_buffers.bind_group_pair.layout, &s.material_buffer.bindgroup_pair.as_ref().unwrap().layout,
|
||||
&s.bgl_texture,
|
||||
&s.light_cull_compute.light_indices_grid.bg_pair.layout,
|
||||
])));
|
||||
&s.light_buffers.bindgroup_layout, &s.material_buffer.bindgroup_pair.as_ref().unwrap().layout,
|
||||
&s.bgl_texture])));
|
||||
s.render_pipelines = pipelines;
|
||||
|
||||
s
|
||||
|
|
@ -351,24 +347,18 @@ impl BasicRenderer {
|
|||
let (vertex_buffer, buffer_indices) = self.create_vertex_index_buffers(mesh);
|
||||
|
||||
let material = mesh.material.as_ref()
|
||||
.expect("Material resource not loaded yet");
|
||||
let material_ref = material.data_ref()
|
||||
.expect("Material resource not loaded yet")
|
||||
.data_ref()
|
||||
.unwrap();
|
||||
|
||||
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);
|
||||
let material = Material::from_resource(&self.device, &self.queue, self.bgl_texture.clone(), &material);
|
||||
let uni = MaterialUniform::from(&material);
|
||||
self.queue.write_buffer(&self.material_buffer.inner_buf, 0, bytemuck::bytes_of(&uni));
|
||||
debug!("Wrote material to buffer");
|
||||
|
||||
MeshBufferStorage {
|
||||
buffer_vertex: vertex_buffer,
|
||||
buffer_indices,
|
||||
material: Some(material.clone()),
|
||||
material: Some(material),
|
||||
}
|
||||
}
|
||||
|
||||
|
|
@ -519,7 +509,7 @@ impl Renderer for BasicRenderer {
|
|||
let output = self.surface.get_current_texture()?;
|
||||
let view = output.texture.create_view(&wgpu::TextureViewDescriptor::default());
|
||||
|
||||
self.light_cull_compute.compute(&self.camera_buffer, &self.light_buffers, &self.depth_buffer_texture);
|
||||
//self.light_cull_compute.compute();
|
||||
|
||||
let mut encoder = self.device.create_command_encoder(&wgpu::CommandEncoderDescriptor {
|
||||
label: Some("Basic Renderer's Encoder")
|
||||
|
|
@ -581,11 +571,9 @@ impl Renderer for BasicRenderer {
|
|||
render_pass.set_bind_group(1, bindgroup, &[ offset, offset, ]);
|
||||
|
||||
render_pass.set_bind_group(2, &self.camera_buffer.bindgroup(), &[]);
|
||||
render_pass.set_bind_group(3, &self.light_buffers.bind_group_pair.bindgroup, &[]);
|
||||
render_pass.set_bind_group(3, &self.light_buffers.bindgroup, &[]);
|
||||
render_pass.set_bind_group(4, &self.material_buffer.bindgroup_pair.as_ref().unwrap().bindgroup, &[]);
|
||||
|
||||
render_pass.set_bind_group(6, &self.light_cull_compute.light_indices_grid.bg_pair.bindgroup, &[]);
|
||||
|
||||
// if this mesh uses indices, use them to draw the mesh
|
||||
if let Some((idx_type, indices)) = buffers.buffer_indices.as_ref() {
|
||||
let indices_len = indices.count() as u32;
|
||||
|
|
@ -617,13 +605,7 @@ impl Renderer for BasicRenderer {
|
|||
|
||||
// tell other things of updated resize
|
||||
self.surface.configure(&self.device, &self.config);
|
||||
|
||||
let create_bindgroup = self.depth_buffer_texture.bindgroup_pair.is_some();
|
||||
self.depth_buffer_texture = RenderTexture::create_depth_texture(&self.device, &self.config, "Depth Buffer Texture");
|
||||
if create_bindgroup {
|
||||
self.depth_buffer_texture.create_bind_group(&self.device);
|
||||
}
|
||||
|
||||
self.inuse_camera.update_aspect_ratio(self.size);
|
||||
self.light_cull_compute.update_screen_size(new_size);
|
||||
}
|
||||
|
|
|
|||
|
|
@ -2,9 +2,9 @@
|
|||
|
||||
const max_light_count: u32 = 16u;
|
||||
|
||||
const LIGHT_TY_DIRECTIONAL = 0u;
|
||||
const LIGHT_TY_POINT = 1u;
|
||||
const LIGHT_TY_SPOT = 2u;
|
||||
const light_ty_directional = 0u;
|
||||
const light_ty_point = 1u;
|
||||
const light_ty_spot = 2u;
|
||||
|
||||
struct VertexInput {
|
||||
@location(0) position: vec3<f32>,
|
||||
|
|
@ -97,79 +97,41 @@ var t_specular: texture_2d<f32>;
|
|||
@group(5) @binding(1)
|
||||
var s_specular: sampler;
|
||||
|
||||
@group(6) @binding(0)
|
||||
var<storage, read_write> u_light_indices: array<u32>;
|
||||
@group(6) @binding(1)
|
||||
var t_light_grid: texture_storage_2d<rg32uint, read_write>; // vec2<u32>
|
||||
|
||||
@fragment
|
||||
fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
|
||||
/*let object_color: vec4<f32> = textureSample(t_diffuse, s_diffuse, in.tex_coords);
|
||||
let object_color: vec4<f32> = textureSample(t_diffuse, s_diffuse, in.tex_coords);
|
||||
let specular_color: vec3<f32> = textureSample(t_specular, s_specular, in.tex_coords).xyz;
|
||||
|
||||
// this needs to be 0.0 for the math
|
||||
//u_lights.directional_light.direction.w = 0.0;
|
||||
|
||||
var light_res = vec3<f32>(0.0);
|
||||
|
||||
for (var i = 0u; i < u_lights.light_count; i++) {
|
||||
var light = u_lights.data[i];
|
||||
|
||||
if (light.light_ty == LIGHT_TY_DIRECTIONAL) {
|
||||
if (light.light_ty == light_ty_directional) {
|
||||
light_res += blinn_phong_dir_light(in.world_position, in.world_normal, light, u_material, specular_color);
|
||||
} else if (light.light_ty == LIGHT_TY_POINT) {
|
||||
} else if (light.light_ty == light_ty_point) {
|
||||
light_res += blinn_phong_point_light(in.world_position, in.world_normal, light, u_material, specular_color);
|
||||
} else if (light.light_ty == LIGHT_TY_SPOT) {
|
||||
} else if (light.light_ty == light_ty_spot) {
|
||||
light_res += blinn_phong_spot_light(in.world_position, in.world_normal, light, u_material, specular_color);
|
||||
}
|
||||
}
|
||||
|
||||
/*var light_res = blinn_phong_dir_light(in.world_position, in.world_normal, u_lights.directional_light, u_material, specular_color);
|
||||
|
||||
for (var i = 0u; i < u_lights.point_light_count; i++) {
|
||||
light_res += blinn_phong_point_light(in.world_position, in.world_normal, u_lights.point_lights[i], u_material, specular_color);
|
||||
}
|
||||
|
||||
for (var i = 0u; i < u_lights.spot_light_count; i++) {
|
||||
light_res += blinn_phong_spot_light(in.world_position, in.world_normal, u_lights.spot_lights[i], u_material, specular_color);
|
||||
}*/
|
||||
|
||||
let light_object_res = light_res * (object_color.xyz/* * u_material.diffuse.xyz*/);
|
||||
|
||||
return vec4<f32>(light_object_res, object_color.a);*/
|
||||
|
||||
let tile_index = vec2<u32>(floor(in.clip_position.xy / 16.0));
|
||||
let tile: vec2<u32> = textureLoad(t_light_grid, tile_index).xy;
|
||||
|
||||
let object_color: vec4<f32> = textureSample(t_diffuse, s_diffuse, in.tex_coords);
|
||||
let specular_color: vec3<f32> = textureSample(t_specular, s_specular, in.tex_coords).xyz;
|
||||
var light_res = vec3<f32>(0.0);
|
||||
|
||||
let light_offset = tile.x;
|
||||
let light_count = tile.y;
|
||||
|
||||
for (var i = 0u; i < light_count; i++) {
|
||||
let light_index = u_light_indices[light_offset + i];
|
||||
let light: Light = u_lights.data[light_index];
|
||||
|
||||
if (light.light_ty == LIGHT_TY_DIRECTIONAL) {
|
||||
light_res += blinn_phong_dir_light(in.world_position, in.world_normal, light, u_material, specular_color);
|
||||
} else if (light.light_ty == LIGHT_TY_POINT) {
|
||||
light_res += blinn_phong_point_light(in.world_position, in.world_normal, light, u_material, specular_color);
|
||||
} else if (light.light_ty == LIGHT_TY_SPOT) {
|
||||
light_res += blinn_phong_spot_light(in.world_position, in.world_normal, light, u_material, specular_color);
|
||||
}
|
||||
}
|
||||
|
||||
let light_object_res = light_res * (object_color.xyz);
|
||||
return vec4<f32>(light_object_res, object_color.a);
|
||||
|
||||
//return debug_grid(in);
|
||||
}
|
||||
|
||||
fn debug_grid(in: VertexOutput) -> vec4<f32> {
|
||||
let tile_index_float: vec2<f32> = in.clip_position.xy / 16.0;
|
||||
let tile_index = vec2<u32>(floor(tile_index_float));
|
||||
let tile: vec2<u32> = textureLoad(t_light_grid, tile_index).xy;
|
||||
|
||||
// detect where the line grids would be at
|
||||
let x = tile_index_float.x - trunc(tile_index_float.x);
|
||||
let y = tile_index_float.y - trunc(tile_index_float.y);
|
||||
let ta: bool = x < 0.05 || y < 0.05;
|
||||
let tb: bool = x > 0.95 || y > 0.95;
|
||||
|
||||
if ( ta || tb ) {
|
||||
return vec4<f32>(0.0, 0.0, 0.0, 1.0);
|
||||
} else {
|
||||
return vec4<f32>(f32(tile_index.x) / 50.0, f32(tile_index.y) / 38.0, 0.0, 1.0);
|
||||
}
|
||||
}
|
||||
|
||||
fn blinn_phong_dir_light(world_pos: vec3<f32>, world_norm: vec3<f32>, dir_light: Light, material: Material, specular_factor: vec3<f32>) -> vec3<f32> {
|
||||
|
|
|
|||
|
|
@ -1,9 +1,9 @@
|
|||
const BLOCK_SIZE: i32 = 16;
|
||||
const MAX_TILE_VISIBLE_LIGHTS: u32 = 1024u;
|
||||
const block_size: i32 = 16;
|
||||
const max_tile_visible_lights: u32 = 1024u;
|
||||
|
||||
const LIGHT_TY_DIRECTIONAL = 0u;
|
||||
const LIGHT_TY_POINT = 1u;
|
||||
const LIGHT_TY_SPOT = 2u;
|
||||
const light_ty_directional = 0u;
|
||||
const light_ty_point = 1u;
|
||||
const light_ty_spot = 2u;
|
||||
|
||||
// Possible computer shader inputs:
|
||||
//
|
||||
|
|
@ -44,13 +44,13 @@ var<workgroup> wg_max_depth: atomic<u32>;
|
|||
var<workgroup> wg_frustum_planes: array<vec4<f32>, 6>;
|
||||
|
||||
// index list of visible light sources for this tile
|
||||
var<workgroup> wg_visible_light_indices: array<u32, MAX_TILE_VISIBLE_LIGHTS>;
|
||||
var<workgroup> wg_visible_light_indices: array<u32, max_tile_visible_lights>;
|
||||
var<workgroup> wg_visible_light_count: atomic<u32>;
|
||||
|
||||
//var<workgroup> view_projection: mat4x4;
|
||||
|
||||
@group(0) @binding(0)
|
||||
var t_depthmap: texture_depth_2d;
|
||||
var t_depthmap: texture_2d<f32>;
|
||||
@group(0) @binding(1)
|
||||
var s_depthmap: sampler;
|
||||
|
||||
|
|
@ -62,8 +62,11 @@ var<storage, read> u_lights: Lights;
|
|||
|
||||
@group(3) @binding(0)
|
||||
var<storage, read_write> u_light_indices: array<u32>;
|
||||
/*@group(3) @binding(1)
|
||||
var<uniform> u_light_grid: array<array<vec2<u32>>>;*/
|
||||
|
||||
@group(3) @binding(1)
|
||||
var t_light_grid: texture_storage_2d<rg32uint, read_write>;
|
||||
var t_light_grid: texture_storage_2d<rg32uint, read_write>; // rg32uint = vec2<u32> or vec4<u32>(r, g, 0.0, 1.0)
|
||||
|
||||
@group(4) @binding(0)
|
||||
var<uniform> u_screen_size: vec2<u32>;
|
||||
|
|
@ -93,8 +96,9 @@ fn cs_main(
|
|||
workgroupBarrier();
|
||||
|
||||
// step 1: calculate the minimum and maximum depth values for this tile (using the depth map)
|
||||
var tex_coord = vec2<u32>(global_invocation_id.xy);
|
||||
var depth_float: f32 = textureLoad(t_depthmap, tex_coord, 0);
|
||||
var tex_coord = vec2<f32>(global_invocation_id.xy);
|
||||
//var depth_float: f32 = textureSample(t_depthmap, s_depthmap, tex_coord).r;
|
||||
var depth_float = 0.0;
|
||||
// bitcast the floating depth to u32 for atomic comparisons between threads
|
||||
var depth_uint: u32 = bitcast<u32>(depth_float);
|
||||
|
||||
|
|
@ -142,7 +146,7 @@ fn cs_main(
|
|||
// Process the lights detecting which ones to cull for this tile.
|
||||
// Processes 256 lights simultaniously, each on a thread in the workgroup. Requires multiple
|
||||
// iterations for more lights.
|
||||
var thread_count = u32(BLOCK_SIZE * BLOCK_SIZE);
|
||||
var thread_count = u32(block_size * block_size);
|
||||
var pass_count = (u_lights.light_count + thread_count - 1u) / thread_count;
|
||||
for (var i = 0u; i < pass_count; i++) {
|
||||
// find the light index to check on this thread, make sure we're not trying to test
|
||||
|
|
@ -156,21 +160,18 @@ fn cs_main(
|
|||
var position = light.position;
|
||||
var radius = light.range;
|
||||
|
||||
if (light.light_ty == LIGHT_TY_DIRECTIONAL) {
|
||||
add_light(light_index);
|
||||
} else if (light.light_ty == LIGHT_TY_POINT
|
||||
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])) {
|
||||
add_light(light_index);
|
||||
/*var offset: u32 = wg_visible_light_count;
|
||||
var offset: u32 = wg_visible_light_count;
|
||||
|
||||
if (offset < MAX_TILE_VISIBLE_LIGHTS) {
|
||||
if (offset < max_tile_visible_lights) {
|
||||
atomicAdd(&wg_visible_light_count, 1u);
|
||||
wg_visible_light_indices[offset] = light_index;
|
||||
}*/
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
@ -181,14 +182,28 @@ fn cs_main(
|
|||
|
||||
// first update the light grid on the first thread
|
||||
if (local_invocation_index == 0u) {
|
||||
var offset = u32(index) * MAX_TILE_VISIBLE_LIGHTS; // index in the global light list
|
||||
var offset = u32(index) * max_tile_visible_lights; // index in the global light list
|
||||
//t_light_grid[workgroup_id.x][workgroup_id.y] = vec2<f32>(offset, wg_visible_light_count);
|
||||
textureStore(t_light_grid, workgroup_id.xy, vec4<u32>(offset, wg_visible_light_count, 0u, 1u));
|
||||
|
||||
// TODO: update transparent light grid
|
||||
|
||||
/*var offset = index * max_tile_visible_lights; // position in the global light buffer
|
||||
// update the light
|
||||
for (var i = 0u; i < wg_visible_light_count; i++) {
|
||||
//u_visible_light_indices[offset + i] = wg_visible_light_indices[i];
|
||||
}
|
||||
|
||||
if (wg_visible_light_count != 1024) {
|
||||
// Mark the end of the visible lights for this tile
|
||||
u_visible_light_indices[offset + wg_visible_light_count] = -1;
|
||||
}*/
|
||||
}
|
||||
|
||||
workgroupBarrier();
|
||||
|
||||
// now update the light index list on all threads.
|
||||
var indices_offset = u32(index) * MAX_TILE_VISIBLE_LIGHTS;
|
||||
var indices_offset = u32(index) * max_tile_visible_lights;
|
||||
//var pass_count = (wg_visible_light_count + thread_count - 1) / thread_count;
|
||||
for (var i = 0u; i < pass_count; i++) {
|
||||
// find the light index to check on this thread, make sure we're not trying to test
|
||||
|
|
@ -209,20 +224,6 @@ fn cs_main(
|
|||
}
|
||||
}
|
||||
|
||||
/// Add a light to the visible light indicies list.
|
||||
/// Returns a boolean indicating if the light was added.
|
||||
fn add_light(light_index: u32) -> bool {
|
||||
var offset: u32 = wg_visible_light_count;
|
||||
|
||||
if (offset < MAX_TILE_VISIBLE_LIGHTS) {
|
||||
atomicAdd(&wg_visible_light_count, 1u);
|
||||
wg_visible_light_indices[offset] = light_index;
|
||||
return true;
|
||||
}
|
||||
|
||||
return false;
|
||||
}
|
||||
|
||||
fn sphere_inside_frustrum(frustum: array<vec4<f32>, 6>, sphere_origin: vec3<f32>, radius: f32) -> bool {
|
||||
// to be able to index this array with a non-const value,
|
||||
// it must be defined as a var
|
||||
|
|
@ -243,7 +244,5 @@ fn sphere_inside_frustrum(frustum: array<vec4<f32>, 6>, sphere_origin: vec3<f32>
|
|||
/// Source: Real-time collision detection, Christer Ericson (2005)
|
||||
/// (https://www.3dgep.com/forward-plus/#light-culling-compute-shader)
|
||||
fn sphere_inside_plane(sphere_origin: vec3<f32>, radius: f32, plane: vec4<f32>) -> bool {
|
||||
//return dot(plane.xyz, sphere_origin) - plane.w < -radius;
|
||||
|
||||
return dot(vec4<f32>(sphere_origin, 0.0), plane) + radius > 0.0;
|
||||
return dot(plane.xyz, sphere_origin) - plane.w < -radius;
|
||||
}
|
||||
|
|
@ -1,4 +1,4 @@
|
|||
use std::rc::Rc;
|
||||
use std::sync::Arc;
|
||||
|
||||
use image::GenericImageView;
|
||||
use lyra_resource::{FilterMode, ResHandle, Texture, WrappingMode};
|
||||
|
|
@ -44,7 +44,7 @@ impl RenderTexture {
|
|||
})
|
||||
}
|
||||
|
||||
fn create_bind_group_pair(device: &wgpu::Device, layout: Rc<wgpu::BindGroupLayout>, view: &wgpu::TextureView, sampler: &wgpu::Sampler) -> BindGroupPair {
|
||||
fn create_bind_group_pair(device: &wgpu::Device, layout: Arc<wgpu::BindGroupLayout>, view: &wgpu::TextureView, sampler: &wgpu::Sampler) -> BindGroupPair {
|
||||
let bg = device.create_bind_group(
|
||||
&wgpu::BindGroupDescriptor {
|
||||
layout: &layout,
|
||||
|
|
@ -68,12 +68,12 @@ impl RenderTexture {
|
|||
}
|
||||
}
|
||||
|
||||
pub fn from_bytes(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Rc<wgpu::BindGroupLayout>, bytes: &[u8], label: &str) -> anyhow::Result<Self> {
|
||||
pub fn from_bytes(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Arc<wgpu::BindGroupLayout>, bytes: &[u8], label: &str) -> anyhow::Result<Self> {
|
||||
let img = image::load_from_memory(bytes)?;
|
||||
Self::from_image(device, queue, bg_layout, &img, Some(label))
|
||||
}
|
||||
|
||||
pub fn from_image(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Rc<wgpu::BindGroupLayout>, img: &image::DynamicImage, label: Option<&str>) -> anyhow::Result<Self> {
|
||||
pub fn from_image(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Arc<wgpu::BindGroupLayout>, img: &image::DynamicImage, label: Option<&str>) -> anyhow::Result<Self> {
|
||||
let rgba = img.to_rgba8();
|
||||
let dimensions = img.dimensions();
|
||||
|
||||
|
|
@ -134,7 +134,7 @@ impl RenderTexture {
|
|||
})
|
||||
}
|
||||
|
||||
pub fn from_resource(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Rc<wgpu::BindGroupLayout>, texture_res: &ResHandle<Texture>, label: Option<&str>) -> anyhow::Result<Self> {
|
||||
pub fn from_resource(device: &wgpu::Device, queue: &wgpu::Queue, bg_layout: Arc<wgpu::BindGroupLayout>, texture_res: &ResHandle<Texture>, label: Option<&str>) -> anyhow::Result<Self> {
|
||||
let texture_ref = texture_res.data_ref().unwrap();
|
||||
let img = texture_ref.image.data_ref().unwrap();
|
||||
|
||||
|
|
|
|||
Loading…
Reference in New Issue