diff --git a/src/client/mod.rs b/src/client/mod.rs index 3a7d9b8..78bf819 100644 --- a/src/client/mod.rs +++ b/src/client/mod.rs @@ -19,7 +19,7 @@ use system::render::add_grid; use crate::{ common::{ClientMessage, ServerHandle, ServerMessage}, - server::Server, + server::Server, util::timer::Timer, }; use self::{input::Input, render::Renderer, ClientState}; @@ -37,6 +37,7 @@ pub struct Client<'a> { window: Arc, state: ClientState, renderer: Renderer<'a>, + render_timer: Timer, render_commands: Vec, exit: bool, input: Input, @@ -78,6 +79,7 @@ impl Client<'_> { window, exit: false, renderer, + render_timer: Timer::new(60), render_commands: Vec::new(), state, input: Input::new(), @@ -132,17 +134,18 @@ impl Client<'_> { commands.extend(world_cmds); self.renderer.handle_commands(commands); self.renderer.draw(); + self.render_timer.add(self.renderer.timer().duration(0)); } if now >= self.second_target { self.second_target += Duration::from_secs(1); - // let timer = self.renderer.timer(); - // println!( - // "avg: {:4?}; max: {:4?}; fps: {:4?}", - // timer.avg(), - // timer.max(), - // timer.per_sec(), - // ); + let timer = &self.render_timer; + println!( + "avg: {:4?}; max: {:4?}; fps: {:4?}", + timer.avg(), + timer.max(), + timer.per_sec(), + ); } if self.exit { diff --git a/src/client/render/command.rs b/src/client/render/command.rs index 3ed91ae..b83e531 100644 --- a/src/client/render/command.rs +++ b/src/client/render/command.rs @@ -7,8 +7,6 @@ use super::{voxel::VoxelColor, Renderer}; use bevy_ecs::entity::Entity; use nalgebra::{Rotation3, Vector3}; use ndarray::Array3; -use std::sync::Arc; -use winit::window::Window; #[derive(Debug, Clone)] pub enum RenderCommand { @@ -76,7 +74,6 @@ impl<'a> Renderer<'a> { &self.device, &mut self.encoder, &mut self.staging_belt, - self.size, &self.camera, ); } diff --git a/src/client/render/mod.rs b/src/client/render/mod.rs index 7f05177..df1e797 100644 --- a/src/client/render/mod.rs +++ b/src/client/render/mod.rs @@ -8,7 +8,7 @@ pub use command::*; use super::camera::Camera; use crate::client::rsc::CLEAR_COLOR; use nalgebra::Vector2; -use util::DepthTexture; +use util::GPUTimer; use voxel::VoxelPipeline; use winit::{dpi::PhysicalSize, window::Window}; @@ -21,14 +21,12 @@ pub struct Renderer<'a> { config: wgpu::SurfaceConfiguration, staging_belt: wgpu::util::StagingBelt, voxel_pipeline: VoxelPipeline, + timer: GPUTimer, camera: Camera, - depth_texture: DepthTexture, } impl<'a> Renderer<'a> { - pub fn new( - window: Arc, - ) -> Self { + pub fn new(window: Arc) -> Self { let size = window.inner_size(); let instance = wgpu::Instance::new(wgpu::InstanceDescriptor { @@ -47,11 +45,13 @@ impl<'a> Renderer<'a> { })) .expect("Could not get adapter!"); - let buf_size = (10u32.pow(9) * 15) / 10; + let buf_size = (10f32.powi(9) * 1.5) as u32; let (device, queue) = pollster::block_on(adapter.request_device( &wgpu::DeviceDescriptor { label: None, - required_features: wgpu::Features::empty(), + required_features: wgpu::Features::TIMESTAMP_QUERY + | wgpu::Features::TIMESTAMP_QUERY_INSIDE_ENCODERS + | wgpu::Features::TIMESTAMP_QUERY_INSIDE_PASSES, required_limits: wgpu::Limits { max_storage_buffer_binding_size: buf_size, max_buffer_size: buf_size as u64, @@ -94,8 +94,6 @@ impl<'a> Renderer<'a> { // doesn't affect performance much and depends on "normal" zoom let staging_belt = wgpu::util::StagingBelt::new(4096 * 4); - let depth_texture = DepthTexture::init(&device, &config, "depth_texture"); - Self { camera: Camera::default(), size: Vector2::new(size.width, size.height), @@ -103,10 +101,10 @@ impl<'a> Renderer<'a> { staging_belt, surface, encoder: Self::create_encoder(&device), + timer: GPUTimer::new(&device, 1), device, config, queue, - depth_texture, } } @@ -135,7 +133,9 @@ impl<'a> Renderer<'a> { label: None, timestamp_writes: None, }); + self.timer.start_compute(&mut compute_pass, 0); self.voxel_pipeline.compute(&mut compute_pass); + self.timer.stop_compute(&mut compute_pass, 0); drop(compute_pass); let mut render_pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { @@ -148,14 +148,6 @@ impl<'a> Renderer<'a> { store: wgpu::StoreOp::Store, }, })], - // depth_stencil_attachment: Some(wgpu::RenderPassDepthStencilAttachment { - // view: &self.depth_texture.view, - // depth_ops: Some(wgpu::Operations { - // load: wgpu::LoadOp::Clear(1.0), - // store: wgpu::StoreOp::Store, - // }), - // stencil_ops: None, - // }), depth_stencil_attachment: None, timestamp_writes: None, occlusion_query_set: None, @@ -163,10 +155,14 @@ impl<'a> Renderer<'a> { self.voxel_pipeline.draw(&mut render_pass); drop(render_pass); + self.timer.resolve(&mut encoder); + self.staging_belt.finish(); self.queue.submit(std::iter::once(encoder.finish())); output.present(); self.staging_belt.recall(); + + self.timer.finish(&self.device); } pub fn resize(&mut self, size: PhysicalSize) { @@ -175,14 +171,15 @@ impl<'a> Renderer<'a> { self.config.height = size.height; self.surface.configure(&self.device, &self.config); self.voxel_pipeline.resize(&self.device, self.size); - - self.depth_texture = DepthTexture::init(&self.device, &self.config, "depth_texture"); self.voxel_pipeline.update_view( &self.device, &mut self.encoder, &mut self.staging_belt, - self.size, &self.camera, ); } + + pub fn timer(&self) -> &GPUTimer { + &self.timer + } } diff --git a/src/client/render/util/buf.rs b/src/client/render/util/array.rs similarity index 72% rename from src/client/render/util/buf.rs rename to src/client/render/util/array.rs index cb253b5..4a4d545 100644 --- a/src/client/render/util/buf.rs +++ b/src/client/render/util/array.rs @@ -1,7 +1,7 @@ use std::marker::PhantomData; use wgpu::{util::DeviceExt, BufferAddress, BufferUsages}; -pub struct ArrBuf { +pub struct ArrayBuffer { len: usize, buffer: wgpu::Buffer, label: String, @@ -10,7 +10,7 @@ pub struct ArrBuf { moves: Vec, } -impl ArrBuf { +impl ArrayBuffer { pub fn update( &mut self, device: &wgpu::Device, @@ -62,6 +62,42 @@ impl ArrBuf { resized } + pub fn add( + &mut self, + device: &wgpu::Device, + encoder: &mut wgpu::CommandEncoder, + belt: &mut wgpu::util::StagingBelt, + data: &[T], + ) { + self.update( + device, + encoder, + belt, + self.len + data.len(), + &[ArrBufUpdate { + offset: self.len, + data, + }], + ); + } + + pub fn set( + &mut self, + device: &wgpu::Device, + encoder: &mut wgpu::CommandEncoder, + belt: &mut wgpu::util::StagingBelt, + offset: usize, + data: &[T], + ) { + self.update( + device, + encoder, + belt, + self.len, + &[ArrBufUpdate { offset, data }], + ); + } + pub fn init(device: &wgpu::Device, label: &str, usage: BufferUsages) -> Self { let label = &(label.to_owned() + " Buffer"); Self { @@ -116,10 +152,6 @@ impl ArrBuf { }) } - pub fn buffer(&self) -> &wgpu::Buffer { - &self.buffer - } - pub fn mov(&mut self, mov: BufMove) { self.moves.push(mov); } @@ -127,6 +159,31 @@ impl ArrBuf { pub fn len(&self) -> usize { self.len } + + pub fn bind_group_layout_entry( + &self, + binding: u32, + visibility: wgpu::ShaderStages, + ty: wgpu::BufferBindingType, + ) -> wgpu::BindGroupLayoutEntry { + wgpu::BindGroupLayoutEntry { + binding, + visibility, + ty: wgpu::BindingType::Buffer { + ty, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + } + } + + pub fn bind_group_entry(&self, binding: u32) -> wgpu::BindGroupEntry { + return wgpu::BindGroupEntry { + binding, + resource: self.buffer.as_entire_binding(), + }; + } } pub struct ArrBufUpdate<'a, T> { diff --git a/src/client/render/util/instance.rs b/src/client/render/util/instance.rs deleted file mode 100644 index 12904d4..0000000 --- a/src/client/render/util/instance.rs +++ /dev/null @@ -1,78 +0,0 @@ -use wgpu::{BufferUsages, VertexAttribute}; - -use super::buf::{ArrBuf, ArrBufUpdate, BufMove}; - -pub struct Instances { - buf: ArrBuf, - location: u32, - attrs: Vec, -} - -impl Instances { - pub fn update( - &mut self, - device: &wgpu::Device, - encoder: &mut wgpu::CommandEncoder, - belt: &mut wgpu::util::StagingBelt, - size: usize, - updates: &[ArrBufUpdate], - ) -> bool { - self.buf.update(device, encoder, belt, size, updates) - } - - pub fn init( - device: &wgpu::Device, - label: &str, - location: u32, - attrs: &[wgpu::VertexAttribute], - ) -> Self { - Self { - buf: ArrBuf::init( - device, - &(label.to_owned() + " Instance"), - BufferUsages::VERTEX, - ), - location, - attrs: attrs.to_vec(), - } - } - - pub fn init_with( - device: &wgpu::Device, - label: &str, - location: u32, - attrs: &[wgpu::VertexAttribute], - data: &[T], - ) -> Self { - Self { - buf: ArrBuf::init_with( - device, - &(label.to_owned() + " Instance"), - BufferUsages::VERTEX, - data, - ), - location, - attrs: attrs.to_vec(), - } - } - - pub fn set_in<'a>(&'a self, render_pass: &mut wgpu::RenderPass<'a>) { - render_pass.set_vertex_buffer(self.location, self.buf.buffer().slice(..)); - } - - pub fn desc(&self) -> wgpu::VertexBufferLayout { - wgpu::VertexBufferLayout { - array_stride: std::mem::size_of::() as wgpu::BufferAddress, - step_mode: wgpu::VertexStepMode::Instance, - attributes: &self.attrs, - } - } - - pub fn mov(&mut self, mov: BufMove) { - self.buf.mov(mov); - } - - pub fn len(&self) -> usize { - self.buf.len() - } -} diff --git a/src/client/render/util/mod.rs b/src/client/render/util/mod.rs index cc11480..794e99e 100644 --- a/src/client/render/util/mod.rs +++ b/src/client/render/util/mod.rs @@ -1,13 +1,9 @@ -mod buf; -mod instance; -mod storage; +mod array; mod uniform; -mod vertex; mod texture; +mod timer; -pub use buf::*; -pub use instance::*; -pub use storage::*; +pub use array::*; pub use uniform::*; -pub use vertex::*; pub use texture::*; +pub use timer::*; diff --git a/src/client/render/util/storage.rs b/src/client/render/util/storage.rs deleted file mode 100644 index a18537a..0000000 --- a/src/client/render/util/storage.rs +++ /dev/null @@ -1,73 +0,0 @@ -use super::buf::{ArrBuf, ArrBufUpdate, BufMove}; -use wgpu::BufferUsages; - -pub struct Storage { - binding: u32, - buf: ArrBuf, - visibility: wgpu::ShaderStages, -} - -impl Storage { - pub fn init(device: &wgpu::Device, visibility: wgpu::ShaderStages, label: &str, binding: u32) -> Self { - Self { - buf: ArrBuf::init( - device, - &(label.to_owned() + " Storage"), - BufferUsages::STORAGE, - ), - visibility, - binding, - } - } - pub fn init_with(device: &wgpu::Device, visibility: wgpu::ShaderStages, label: &str, binding: u32, data: &[T]) -> Self { - Self { - buf: ArrBuf::init_with( - device, - &(label.to_owned() + " Storage"), - BufferUsages::STORAGE, - data - ), - visibility, - binding, - } - } -} - -impl Storage { - pub fn bind_group_layout_entry(&self) -> wgpu::BindGroupLayoutEntry { - wgpu::BindGroupLayoutEntry { - binding: self.binding, - visibility: self.visibility, - ty: wgpu::BindingType::Buffer { - ty: wgpu::BufferBindingType::Storage { read_only: true }, - has_dynamic_offset: false, - min_binding_size: None, - }, - count: None, - } - } - pub fn bind_group_entry(&self) -> wgpu::BindGroupEntry { - return wgpu::BindGroupEntry { - binding: self.binding, - resource: self.buf.buffer().as_entire_binding(), - }; - } - pub fn update( - &mut self, - device: &wgpu::Device, - encoder: &mut wgpu::CommandEncoder, - belt: &mut wgpu::util::StagingBelt, - size: usize, - updates: &[ArrBufUpdate], - ) -> bool { - self.buf.update(device, encoder, belt, size, updates) - } - - pub fn mov(&mut self, mov: BufMove) { - self.buf.mov(mov); - } - - pub fn len(&mut self) -> usize { - self.buf.len() - } -} diff --git a/src/client/render/util/texture.rs b/src/client/render/util/texture.rs index d626576..49bd121 100644 --- a/src/client/render/util/texture.rs +++ b/src/client/render/util/texture.rs @@ -1,23 +1,26 @@ -pub struct DepthTexture { +pub struct Texture { + texture_desc: wgpu::TextureDescriptor<'static>, + view_desc: wgpu::TextureViewDescriptor<'static>, + sampler_desc: wgpu::SamplerDescriptor<'static>, pub texture: wgpu::Texture, pub view: wgpu::TextureView, pub sampler: wgpu::Sampler, } -impl DepthTexture { +impl Texture { pub const DEPTH_FORMAT: wgpu::TextureFormat = wgpu::TextureFormat::Depth32Float; - pub fn init( + pub fn init_depth( device: &wgpu::Device, config: &wgpu::SurfaceConfiguration, - label: &str, + label: &'static str, ) -> Self { let size = wgpu::Extent3d { width: config.width + 1, height: config.height + 1, depth_or_array_layers: 1, }; - let desc = wgpu::TextureDescriptor { + let texture_desc = wgpu::TextureDescriptor { label: Some(label), size, mip_level_count: 1, @@ -27,97 +30,62 @@ impl DepthTexture { usage: wgpu::TextureUsages::RENDER_ATTACHMENT | wgpu::TextureUsages::TEXTURE_BINDING, view_formats: &[], }; - let texture = device.create_texture(&desc); - - let view = texture.create_view(&wgpu::TextureViewDescriptor::default()); - let sampler = device.create_sampler(&wgpu::SamplerDescriptor { - address_mode_u: wgpu::AddressMode::ClampToEdge, - address_mode_v: wgpu::AddressMode::ClampToEdge, - address_mode_w: wgpu::AddressMode::ClampToEdge, - mag_filter: wgpu::FilterMode::Linear, - min_filter: wgpu::FilterMode::Linear, - mipmap_filter: wgpu::FilterMode::Nearest, - compare: Some(wgpu::CompareFunction::LessEqual), - lod_min_clamp: 0.0, - lod_max_clamp: 100.0, - ..Default::default() - }); + Self::init( + device, + texture_desc, + wgpu::TextureViewDescriptor::default(), + wgpu::SamplerDescriptor { + address_mode_u: wgpu::AddressMode::ClampToEdge, + address_mode_v: wgpu::AddressMode::ClampToEdge, + address_mode_w: wgpu::AddressMode::ClampToEdge, + mag_filter: wgpu::FilterMode::Linear, + min_filter: wgpu::FilterMode::Linear, + mipmap_filter: wgpu::FilterMode::Nearest, + compare: Some(wgpu::CompareFunction::LessEqual), + lod_min_clamp: 0.0, + lod_max_clamp: 100.0, + ..Default::default() + }, + ) + } + pub fn init( + device: &wgpu::Device, + texture_desc: wgpu::TextureDescriptor<'static>, + view_desc: wgpu::TextureViewDescriptor<'static>, + sampler_desc: wgpu::SamplerDescriptor<'static>, + ) -> Self { + let texture = device.create_texture(&texture_desc); + let view = texture.create_view(&view_desc); + let sampler = device.create_sampler(&sampler_desc); Self { + texture_desc, + view_desc, + sampler_desc, texture, view, sampler, } } -} -pub struct StorageTexture { - binding: u32, - visibility: wgpu::ShaderStages, - pub buf: wgpu::Texture, - pub view: wgpu::TextureView, - pub sampler: wgpu::Sampler, -} - -impl StorageTexture { - pub const FORMAT: wgpu::TextureFormat = wgpu::TextureFormat::Rgba8Unorm; - - pub fn init( - device: &wgpu::Device, - size: wgpu::Extent3d, - label: &str, - visibility: wgpu::ShaderStages, - binding: u32, - ) -> Self { - let desc = wgpu::TextureDescriptor { - label: Some(label), - size, - mip_level_count: 1, - sample_count: 1, - dimension: wgpu::TextureDimension::D2, - format: Self::FORMAT, - usage: wgpu::TextureUsages::STORAGE_BINDING | wgpu::TextureUsages::TEXTURE_BINDING, - view_formats: &[], - }; - let texture = device.create_texture(&desc); - - let view = texture.create_view(&wgpu::TextureViewDescriptor::default()); - let sampler = device.create_sampler(&wgpu::SamplerDescriptor { - address_mode_u: wgpu::AddressMode::ClampToEdge, - address_mode_v: wgpu::AddressMode::ClampToEdge, - address_mode_w: wgpu::AddressMode::ClampToEdge, - mag_filter: wgpu::FilterMode::Linear, - min_filter: wgpu::FilterMode::Linear, - mipmap_filter: wgpu::FilterMode::Linear, - compare: None, - lod_min_clamp: 0.0, - lod_max_clamp: 100.0, - ..Default::default() - }); - Self { - visibility, + pub fn resize(&mut self, device: &wgpu::Device, size: wgpu::Extent3d) { + self.texture_desc.size = size; + self.texture = device.create_texture(&self.texture_desc); + self.view = self.texture.create_view(&self.view_desc); + } + pub fn view_bind_group_entry(&self, binding: u32) -> wgpu::BindGroupEntry { + wgpu::BindGroupEntry { binding, - buf: texture, - view, - sampler, - } - } - pub fn bind_group_layout_entry(&self) -> wgpu::BindGroupLayoutEntry { - wgpu::BindGroupLayoutEntry { - binding: self.binding, - visibility: self.visibility, - ty: wgpu::BindingType::StorageTexture { - access: wgpu::StorageTextureAccess::WriteOnly, - format: Self::FORMAT, - view_dimension: wgpu::TextureViewDimension::D2, - }, - count: None, - } - } - pub fn bind_group_entry(&self) -> wgpu::BindGroupEntry { - return wgpu::BindGroupEntry { - binding: self.binding, resource: wgpu::BindingResource::TextureView(&self.view), - }; + } + } + pub fn sampler_bind_group_entry(&self, binding: u32) -> wgpu::BindGroupEntry { + wgpu::BindGroupEntry { + binding, + resource: wgpu::BindingResource::Sampler(&self.sampler), + } + } + pub fn format(&self) -> wgpu::TextureFormat { + self.texture_desc.format } } diff --git a/src/client/render/util/timer.rs b/src/client/render/util/timer.rs new file mode 100644 index 0000000..a9f2c42 --- /dev/null +++ b/src/client/render/util/timer.rs @@ -0,0 +1,85 @@ +use std::time::Duration; + +pub struct GPUTimer { + resolve_buf: wgpu::Buffer, + map_buf: wgpu::Buffer, + query_set: wgpu::QuerySet, + timestamps: Vec, +} + +impl GPUTimer { + pub fn new(device: &wgpu::Device, count: u32) -> Self { + let count = count * 2; + let timestamp_set = device.create_query_set(&wgpu::QuerySetDescriptor { + count, + label: Some("voxel timestamp"), + ty: wgpu::QueryType::Timestamp, + }); + let timestamp_resolve_buf = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("voxel timestamp"), + mapped_at_creation: false, + size: 8 * count as u64, + usage: wgpu::BufferUsages::QUERY_RESOLVE | wgpu::BufferUsages::COPY_SRC, + }); + + let timestamp_mapped_buf = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("voxel timestamp"), + mapped_at_creation: false, + size: 8 * count as u64, + usage: wgpu::BufferUsages::COPY_DST | wgpu::BufferUsages::MAP_READ, + }); + + Self { + query_set: timestamp_set, + resolve_buf: timestamp_resolve_buf, + map_buf: timestamp_mapped_buf, + timestamps: vec![0; count as usize], + } + } + + pub fn resolve(&self, encoder: &mut wgpu::CommandEncoder) { + encoder.resolve_query_set(&self.query_set, 0..2, &self.resolve_buf, 0); + encoder.copy_buffer_to_buffer(&self.resolve_buf, 0, &self.map_buf, 0, self.map_buf.size()); + } + + pub fn duration(&self, i: u32) -> Duration { + let i = i as usize * 2; + Duration::from_nanos(self.timestamps[i + 1] - self.timestamps[i]) + } + + pub fn finish(&mut self, device: &wgpu::Device) { + let (s, r) = std::sync::mpsc::channel(); + self.map_buf + .slice(..) + .map_async(wgpu::MapMode::Read, move |v| { + s.send(v).expect("what"); + }); + device.poll(wgpu::Maintain::wait()).panic_on_timeout(); + if let Ok(Ok(())) = r.recv() { + let data = self.map_buf.slice(..).get_mapped_range(); + self.timestamps.copy_from_slice(bytemuck::cast_slice(&data)); + drop(data); + self.map_buf.unmap(); + } + } + + #[allow(dead_code)] + pub fn start(&self, encoder: &mut wgpu::CommandEncoder, i: u32) { + encoder.write_timestamp(&self.query_set, i * 2); + } + + #[allow(dead_code)] + pub fn stop(&self, encoder: &mut wgpu::CommandEncoder, i: u32) { + encoder.write_timestamp(&self.query_set, i * 2 + 1); + } + + #[allow(dead_code)] + pub fn start_compute(&self, pass: &mut wgpu::ComputePass, i: u32) { + pass.write_timestamp(&self.query_set, i * 2); + } + + #[allow(dead_code)] + pub fn stop_compute(&self, pass: &mut wgpu::ComputePass, i: u32) { + pass.write_timestamp(&self.query_set, i * 2 + 1); + } +} diff --git a/src/client/render/util/vertex.rs b/src/client/render/util/vertex.rs deleted file mode 100644 index 9d892fc..0000000 --- a/src/client/render/util/vertex.rs +++ /dev/null @@ -1,86 +0,0 @@ -use wgpu::{BufferUsages, VertexAttribute}; - -use super::buf::{ArrBuf, ArrBufUpdate, BufMove}; - -pub struct Vertices { - buf: ArrBuf, - location: u32, - attrs: [VertexAttribute; 1], -} - -impl Vertices { - pub fn update( - &mut self, - device: &wgpu::Device, - encoder: &mut wgpu::CommandEncoder, - belt: &mut wgpu::util::StagingBelt, - size: usize, - updates: &[ArrBufUpdate], - ) -> bool { - self.buf.update(device, encoder, belt, size, updates) - } - - pub fn init( - device: &wgpu::Device, - label: &str, - location: u32, - format: wgpu::VertexFormat, - ) -> Self { - Self { - buf: ArrBuf::init( - device, - &(label.to_owned() + " Instance"), - BufferUsages::VERTEX, - ), - location, - attrs: [wgpu::VertexAttribute { - format, - offset: 0, - shader_location: location, - }], - } - } - - pub fn init_with( - device: &wgpu::Device, - label: &str, - location: u32, - format: wgpu::VertexFormat, - data: &[T], - ) -> Self { - Self { - buf: ArrBuf::init_with( - device, - &(label.to_owned() + " Instance"), - BufferUsages::VERTEX, - data, - ), - location, - attrs: [wgpu::VertexAttribute { - format, - offset: 0, - shader_location: location, - }], - } - } - - pub fn set_in<'a>(&'a self, render_pass: &mut wgpu::RenderPass<'a>) { - render_pass.set_vertex_buffer(self.location, self.buf.buffer().slice(..)); - } - - pub fn desc(&self) -> wgpu::VertexBufferLayout { - wgpu::VertexBufferLayout { - array_stride: std::mem::size_of::() as wgpu::BufferAddress, - step_mode: wgpu::VertexStepMode::Vertex, - attributes: &self.attrs, - } - } - - pub fn mov(&mut self, mov: BufMove) { - self.buf.mov(mov); - } - - pub fn len(&self) -> usize { - self.buf.len() - } -} diff --git a/src/client/render/voxel/ray_oct/chunk.rs b/src/client/render/voxel/ray_oct/chunk.rs new file mode 100644 index 0000000..80c5d5b --- /dev/null +++ b/src/client/render/voxel/ray_oct/chunk.rs @@ -0,0 +1,7 @@ +#[repr(C, align(16))] +#[derive(Debug, Clone, Copy, PartialEq, bytemuck::Zeroable)] +pub struct Chunk { + pub offset: u32, +} + +unsafe impl bytemuck::Pod for Chunk {} diff --git a/src/client/render/voxel/ray_oct/group.rs b/src/client/render/voxel/ray_oct/group.rs deleted file mode 100644 index f82cdc8..0000000 --- a/src/client/render/voxel/ray_oct/group.rs +++ /dev/null @@ -1,12 +0,0 @@ -use nalgebra::{Projective3, Vector3}; - -#[repr(C, align(16))] -#[derive(Debug, Clone, Copy, PartialEq, bytemuck::Zeroable)] -pub struct VoxelGroup { - pub transform: Projective3, - pub transform_inv: Projective3, - pub scale: u32, - pub offset: u32, -} - -unsafe impl bytemuck::Pod for VoxelGroup {} diff --git a/src/client/render/voxel/ray_oct/layout.rs b/src/client/render/voxel/ray_oct/layout.rs index 0f96aa7..a382dc0 100644 --- a/src/client/render/voxel/ray_oct/layout.rs +++ b/src/client/render/voxel/ray_oct/layout.rs @@ -1,18 +1,18 @@ -use wgpu::TextureFormat; +use wgpu::{BufferUsages, ShaderStages, TextureFormat}; -use super::{group::VoxelGroup, light::GlobalLight, view::View}; +use super::{chunk::Chunk, light::GlobalLight, view::View}; use crate::{ - client::render::util::{Storage, StorageTexture, Uniform}, + client::render::util::{ArrayBuffer, Texture, Uniform}, util::oct_tree::OctNode, }; use nalgebra::Vector3; pub struct Layout { - pub texture: StorageTexture, pub view: Uniform, - pub voxel_groups: Storage, - pub voxels: Storage, - pub global_lights: Storage, + pub chunks: ArrayBuffer, + pub voxel_data: ArrayBuffer, + pub global_lights: ArrayBuffer, + pub texture: Texture, render_bind_layout: wgpu::BindGroupLayout, compute_bind_layout: wgpu::BindGroupLayout, render_pipeline_layout: wgpu::PipelineLayout, @@ -23,27 +23,40 @@ pub struct Layout { impl Layout { pub fn init(device: &wgpu::Device, config: &wgpu::SurfaceConfiguration) -> Self { let view = Uniform::init(device, "view", 0); - let voxels = Storage::init(device, wgpu::ShaderStages::COMPUTE, "voxels", 1); - let voxel_groups = Storage::init(device, wgpu::ShaderStages::COMPUTE, "voxel groups", 2); - let global_lights = Storage::init_with( + let chunks = ArrayBuffer::init(device, "chunks", BufferUsages::STORAGE); + let voxel_data = ArrayBuffer::init_with( + device, + "voxel data", + BufferUsages::STORAGE, + &[OctNode::new_leaf(0)], + ); + let global_lights = ArrayBuffer::init_with( device, - wgpu::ShaderStages::COMPUTE, "global lights", - 3, + BufferUsages::STORAGE, &[GlobalLight { direction: Vector3::new(-1.0, -2.3, 2.0).normalize(), }], ); - let texture = StorageTexture::init( - device, - wgpu::Extent3d { + let desc = wgpu::TextureDescriptor { + label: Some("compute output"), + size: wgpu::Extent3d { width: config.width, height: config.height, depth_or_array_layers: 1, }, - "compute output", - wgpu::ShaderStages::COMPUTE | wgpu::ShaderStages::FRAGMENT, - 4, + mip_level_count: 1, + sample_count: 1, + dimension: wgpu::TextureDimension::D2, + format: wgpu::TextureFormat::Rgba8Unorm, + usage: wgpu::TextureUsages::STORAGE_BINDING | wgpu::TextureUsages::TEXTURE_BINDING, + view_formats: &[], + }; + let texture = Texture::init( + device, + desc, + wgpu::TextureViewDescriptor::default(), + wgpu::SamplerDescriptor::default(), ); let render_bind_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { @@ -73,10 +86,31 @@ impl Layout { device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { entries: &[ view.bind_group_layout_entry(), - voxels.bind_group_layout_entry(), - voxel_groups.bind_group_layout_entry(), - global_lights.bind_group_layout_entry(), - texture.bind_group_layout_entry(), + chunks.bind_group_layout_entry( + 1, + ShaderStages::COMPUTE, + wgpu::BufferBindingType::Storage { read_only: true }, + ), + voxel_data.bind_group_layout_entry( + 2, + ShaderStages::COMPUTE, + wgpu::BufferBindingType::Storage { read_only: true }, + ), + global_lights.bind_group_layout_entry( + 3, + ShaderStages::COMPUTE, + wgpu::BufferBindingType::Storage { read_only: true }, + ), + wgpu::BindGroupLayoutEntry { + binding: 4, + visibility: ShaderStages::COMPUTE, + ty: wgpu::BindingType::StorageTexture { + access: wgpu::StorageTextureAccess::WriteOnly, + format: texture.format(), + view_dimension: wgpu::TextureViewDimension::D2, + }, + count: None, + }, ], label: Some("voxel compute"), }); @@ -94,8 +128,8 @@ impl Layout { }); Self { view, - voxels, - voxel_groups, + voxel_data, + chunks, global_lights, texture, render_bind_layout, @@ -110,16 +144,10 @@ impl Layout { device.create_bind_group(&wgpu::BindGroupDescriptor { layout: &self.render_bind_layout, entries: &[ - wgpu::BindGroupEntry { - binding: 0, - resource: wgpu::BindingResource::TextureView(&self.texture.view), - }, - wgpu::BindGroupEntry { - binding: 1, - resource: wgpu::BindingResource::Sampler(&self.texture.sampler), - }, + self.texture.view_bind_group_entry(0), + self.texture.sampler_bind_group_entry(1), ], - label: Some("tile_bind_group"), + label: Some("voxel render"), }) } @@ -128,10 +156,10 @@ impl Layout { layout: &self.compute_bind_layout, entries: &[ self.view.bind_group_entry(), - self.voxels.bind_group_entry(), - self.voxel_groups.bind_group_entry(), - self.global_lights.bind_group_entry(), - self.texture.bind_group_entry(), + self.chunks.bind_group_entry(1), + self.voxel_data.bind_group_entry(2), + self.global_lights.bind_group_entry(3), + self.texture.view_bind_group_entry(4), ], label: Some("voxel compute"), }) diff --git a/src/client/render/voxel/ray_oct/mod.rs b/src/client/render/voxel/ray_oct/mod.rs index 1204fd0..180e554 100644 --- a/src/client/render/voxel/ray_oct/mod.rs +++ b/src/client/render/voxel/ray_oct/mod.rs @@ -1,28 +1,25 @@ +mod chunk; mod color; mod grid; -mod group; mod layout; mod light; mod view; use super::super::UpdateGridTransform; -use crate::{ - client::{ - camera::Camera, - render::{ - util::{ArrBufUpdate, StorageTexture}, - AddChunk, CreateVoxelGrid, - }, - }, - common::component::chunk, +use crate::client::{ + camera::Camera, + render::{AddChunk, CreateVoxelGrid}, }; use bevy_ecs::entity::Entity; pub use color::*; use layout::Layout; -use nalgebra::{Projective3, Transform3, Translation3, Vector2, Vector3}; -use std::{collections::HashMap, ops::Deref}; +use nalgebra::{Transform3, Translation3, Vector2}; +use std::{ + collections::HashMap, + sync::{Arc, Mutex}, +}; use wgpu::include_wgsl; -use {group::VoxelGroup, view::View}; +use {chunk::Chunk, view::View}; pub struct VoxelPipeline { layout: Layout, @@ -30,7 +27,7 @@ pub struct VoxelPipeline { compute_bind_group: wgpu::BindGroup, render_pipeline: wgpu::RenderPipeline, render_bind_group: wgpu::BindGroup, - id_map: HashMap, + id_map: HashMap, } const RENDER_SHADER: wgpu::ShaderModuleDescriptor<'_> = include_wgsl!("shader/render.wgsl"); @@ -82,7 +79,8 @@ impl VoxelPipeline { pub fn update_shader(&mut self, device: &wgpu::Device) { let Ok(shader) = std::fs::read_to_string( - env!("CARGO_MANIFEST_DIR").to_owned() + "/src/client/render/voxel/ray_oct/shader/compute.wgsl", + env!("CARGO_MANIFEST_DIR").to_owned() + + "/src/client/render/voxel/ray_oct/shader/compute.wgsl", ) else { println!("Failed to reload shader!"); return; @@ -110,49 +108,27 @@ impl VoxelPipeline { belt: &mut wgpu::util::StagingBelt, AddChunk { id, pos, tree, .. }: AddChunk, ) { - let offset = self.layout.voxels.len(); - - let data = tree.raw(); - let updates = [ArrBufUpdate { offset, data }]; - let size = offset + data.len(); + let offset = self.layout.voxel_data.len() as u32; self.layout - .voxels - .update(device, encoder, belt, size, &updates); + .voxel_data + .add(device, encoder, belt, tree.raw()); - let proj = Projective3::identity() - * Translation3::from((pos.deref() * chunk::SIDE_LENGTH as i32).cast()) - * Translation3::from(-chunk::DIMENSIONS.cast() / 2.0); - let group = VoxelGroup { - transform: proj, - transform_inv: proj.inverse(), - scale: chunk::SCALE, - offset: offset as u32, - }; - let updates = [ArrBufUpdate { - offset: self.layout.voxel_groups.len(), - data: &[group], - }]; - let i = self.layout.voxel_groups.len(); - let size = i + 1; - self.layout - .voxel_groups - .update(device, encoder, belt, size, &updates); + let chunk = Chunk { offset }; + let i = self.layout.chunks.len(); + self.layout.chunks.add(device, encoder, belt, &[chunk]); - self.id_map.insert(id, (i, group)); + self.id_map.insert(id, (i, chunk)); self.compute_bind_group = self.layout.compute_bind_group(device); } pub fn resize(&mut self, device: &wgpu::Device, size: Vector2) { - self.layout.texture = StorageTexture::init( + self.layout.texture.resize( device, wgpu::Extent3d { width: size.x, height: size.y, depth_or_array_layers: 1, }, - "idk man im tired", - wgpu::ShaderStages::COMPUTE | wgpu::ShaderStages::FRAGMENT, - 4, ); self.compute_bind_group = self.layout.compute_bind_group(device); self.render_bind_group = self.layout.render_bind_group(device); @@ -166,21 +142,7 @@ impl VoxelPipeline { update: UpdateGridTransform, ) { if let Some((i, group)) = self.id_map.get_mut(&update.id) { - let offset = Vector3::from_element(-(2u32.pow(group.scale) as f32) / 2.0); - let proj = Projective3::identity() - * Translation3::from(update.pos) - * update.orientation - * Translation3::from(offset); - group.transform = proj; - group.transform_inv = proj.inverse(); - let updates = [ArrBufUpdate { - offset: *i, - data: &[*group], - }]; - let size = self.layout.voxel_groups.len(); - self.layout - .voxel_groups - .update(device, encoder, belt, size, &updates); + self.layout.chunks.set(device, encoder, belt, *i, &[*group]); } } @@ -189,7 +151,6 @@ impl VoxelPipeline { device: &wgpu::Device, encoder: &mut wgpu::CommandEncoder, belt: &mut wgpu::util::StagingBelt, - size: Vector2, camera: &Camera, ) { let transform = @@ -197,6 +158,7 @@ impl VoxelPipeline { let data = View { zoom: camera.scale, transform, + ..Default::default() }; self.layout.view.update(device, encoder, belt, data); } @@ -212,7 +174,7 @@ impl VoxelPipeline { pub fn compute(&self, pass: &mut wgpu::ComputePass) { pass.set_pipeline(&self.compute_pipeline); pass.set_bind_group(0, &self.compute_bind_group, &[]); - let buf = &self.layout.texture.buf; + let buf = &self.layout.texture.texture; let x = (buf.width() - 1) / Self::WORKGROUP_SIZE + 1; let y = (buf.height() - 1) / Self::WORKGROUP_SIZE + 1; pass.dispatch_workgroups(x, y, 1); diff --git a/src/client/render/voxel/ray_oct/shader/compute.wgsl b/src/client/render/voxel/ray_oct/shader/compute.wgsl index 3296af8..4ba174b 100644 --- a/src/client/render/voxel/ray_oct/shader/compute.wgsl +++ b/src/client/render/voxel/ray_oct/shader/compute.wgsl @@ -1,9 +1,9 @@ @group(0) @binding(0) var view: View; @group(0) @binding(1) -var voxels: array; +var chunks: array; @group(0) @binding(2) -var voxel_groups: array; +var voxel_data: array; @group(0) @binding(3) var global_lights: array; @group(0) @binding(4) @@ -16,12 +16,11 @@ struct GlobalLight { struct View { transform: mat4x4, zoom: f32, + chunk_scale: u32, + chunk_dist: u32, }; -struct VoxelGroup { - transform: mat4x4, - transform_inv: mat4x4, - scale: u32, +struct Chunk { offset: u32, }; @@ -38,7 +37,8 @@ fn main(@builtin(global_invocation_id) cell: vec3) { let pixel_pos = vec2( (vec2(cell.xy) / view_dim_f - vec2(0.5)) * vec2(2.0, -2.0 * aspect) ); - let pos = view.transform * vec4(pixel_pos, 1.0, 1.0); + let offset = vec3(f32(1u << (view.chunk_scale - 1))); + let pos = view.transform * vec4(pixel_pos, 1.0, 1.0) + vec4(offset, 0.0); let dir = view.transform * vec4(normalize(vec3(pixel_pos, view.zoom)), 0.0); var color = trace_full(pos, dir); @@ -53,6 +53,7 @@ fn main(@builtin(global_invocation_id) cell: vec3) { const LEAF_BIT = 1u << 31u; const LEAF_MASK = ~LEAF_BIT; +const MAX_HITS = 10; const ZERO3F = vec3(0.0); const ZERO2F = vec2(0.0); @@ -62,126 +63,122 @@ const MAX_ITERS = 10000; // NOTE: CANNOT GO HIGHER THAN 23 due to how floating point // numbers are stored and the bit manipulation used const MAX_SCALE: u32 = 13; -const AMBIENT: f32 = 0.2; -const SPECULAR: f32 = 0.5; - fn trace_full(pos_view: vec4, dir_view: vec4) -> vec4 { - let gi = 0; - let group = voxel_groups[gi]; - if group.scale == 0 { + if arrayLength(&voxel_data) == 1 { return vec4(0.0); } - let dimensions = vec3(1u << group.scale); + let gi = 0; + let chunk = chunks[gi]; + let side_len = 1u << view.chunk_scale; + let dimensions = vec3(side_len); let dim_f = vec3(dimensions); - let dim_i = vec3(dimensions); - // transform so that group is at 0,0 - let pos_start = (group.transform_inv * pos_view).xyz; - var dir = (group.transform_inv * dir_view).xyz; - if dir.x == 0 {dir.x = EPSILON;} - if dir.y == 0 {dir.y = EPSILON;} - if dir.z == 0 {dir.z = EPSILON;} + let pos_start = pos_view.xyz; + var dir = dir_view.xyz; + if dir.x == 0 { dir.x = EPSILON; } + if dir.y == 0 { dir.y = EPSILON; } + if dir.z == 0 { dir.z = EPSILON; } let dir_if = sign(dir); let dir_uf = max(dir_if, vec3(0.0)); + // find where ray intersects with group + // closest (min) and furthest (max) corners of cube relative to direction + let pos_min = (vec3(1.0) - dir_uf) * dim_f; + let pos_max = dir_uf * dim_f; + // time of intersection; x = td + p, solve for t + let t_min = (pos_min - pos_start) / dir; + let t_max = (pos_max - pos_start) / dir; + // time of entrance and exit of the cube + let t_start = max(max(t_min.x, t_min.y), t_min.z); + let t_end = min(min(t_max.x, t_max.y), t_max.z); + if t_end < t_start { return vec4(0.0); } + // axis of intersection + let axis = select(select(2u, 1u, t_start == t_min.y), 0u, t_start == t_min.x); + // time to move entire side length in each direction + let inc_t = abs(1.0 / dir) * f32(side_len); + let t = max(0.0, t_start); + let inv_dir_bits = 7 - vec_to_dir(vec3(dir_uf)); + let corner_adj = t_min - inc_t; // calculate normals var normals = mat3x3( - (group.transform * vec4(dir_if.x, 0.0, 0.0, 0.0)).xyz, - (group.transform * vec4(0.0, dir_if.y, 0.0, 0.0)).xyz, - (group.transform * vec4(0.0, 0.0, dir_if.z, 0.0)).xyz, + vec3(dir_if.x, 0.0, 0.0), + vec3(0.0, dir_if.y, 0.0), + vec3(0.0, 0.0, dir_if.z), ); - var axis = 0u; - // find where ray intersects with group - let pos_min = (vec3(1.0) - dir_uf) * dim_f; - // time of intersection; x = td + p, solve for t - var t_min = (pos_min - pos_start) / dir; - if outside3f(pos_start, ZERO3F, dim_f) { - // points of intersection - let px = pos_start + t_min.x * dir; - let py = pos_start + t_min.y * dir; - let pz = pos_start + t_min.z * dir; + let result = cast_ray(chunk.offset, t, axis, inv_dir_bits, inc_t, corner_adj); + return shade_ray(result, pos_start, dir_view.xyz, t_end, normals); +} - // check if point is in bounds - let hit = vec3( - inside2f(px.yz, ZERO2F, dim_f.yz), - inside2f(py.xz, ZERO2F, dim_f.xz), - inside2f(pz.xy, ZERO2F, dim_f.xy), - ) && (t_min > ZERO3F); - if !any(hit) { - return vec4(0.0); - } - axis = select(select(2u, 1u, hit.y), 0u, hit.x); - } - let t_mult = f32(1u << (MAX_SCALE - group.scale)); - t_min *= t_mult; - // time to move 1 unit in each direction - let full = f32(1u << MAX_SCALE); - let inc_t = abs(1.0 / dir) * full; - let t_offset = max(max(t_min.x, t_min.y), t_min.z); - var t = max(0.0, t_offset); +fn shade_ray(result: RayResult, pos_start: vec3, dir: vec3, t_end: f32, normals: mat3x3) -> vec4 { + var hits = result.hits; - let dir_i = vec3(dir_if); - let dir_u = vec3(dir_uf); - let dir_bits = vec_to_dir(dir_u); - let inv_dir_bits = 7 - dir_bits; - - var node_start = 1u; - var scale = MAX_SCALE - 1; - var scale_exp2 = 0.5; var color = vec4(0.0); - var parents = array(); - var prev = LEAF_BIT; - var old_t = t / t_mult; + for (var i = 0u; i < result.len; i += 1u) { + let hit = hits[i]; + let id = hit.id; + let t = hit.t; + let axis = hit.axis; - var child = 0u; + let next_t = select(hits[i + 1].t, t_end, i == result.len - 1); + + var pos = pos_start + dir * t; + pos[axis] = round(pos[axis]) - f32(dir[axis] < 0.0); + let normal = select(select(normals[0], normals[1], axis == 1), normals[2], axis == 2); + let vcolor = shade(id, pos, normal, dir, next_t - t); + color += vcolor * (1.0 - color.a); + if color.a > FULL_ALPHA { break; } + } + return color; +} + +struct RayHit { + t: f32, + id: u32, + axis: u32, +} + +struct RayResult { + hits: array, + len: u32, +} + +fn cast_ray( + data_offset: u32, t_start: f32, axis_start: u32, + inv_dir_bits: u32, inc_t: vec3, corner_adj: vec3 +) -> RayResult { + var hits = array(); + var depth = 0u; + var min_alpha = 0.0; + + var t = t_start; + var axis = axis_start; + var node_start = 0u; + var scale = MAX_SCALE; + var scale_exp2 = 1.0; + var parents = array(); + var child = inv_dir_bits; var vox_pos = vec3(1.0); - let t_center = t_min + scale_exp2 * inc_t; - if t > t_center.x { vox_pos.x = 1.5; child |= 4u; } - if t > t_center.y { vox_pos.y = 1.5; child |= 2u; } - if t > t_center.z { vox_pos.z = 1.5; child |= 1u; } - let min_adj = t_min - inc_t; + var prev = 0u; var iters = 0; loop { - if iters == MAX_ITERS { - return vec4(1.0, 0.0, 1.0, 1.0); - } + if iters == MAX_ITERS { break; } iters += 1; - let t_corner = vox_pos * inc_t + min_adj; - let node = voxels[group.offset + node_start + (child ^ inv_dir_bits)]; + let t_corner = vox_pos * inc_t + corner_adj; + let node = voxel_data[data_offset + node_start + (child ^ inv_dir_bits)]; if node >= LEAF_BIT { + // ignore consecutive identical leaves if node != prev { - if node != LEAF_BIT { - let real_t = t / t_mult; - let dist = real_t - old_t; - old_t = real_t; - let filt = min(dist / 64.0, 1.0); - if prev == LEAF_BIT + 3 { - color.a += filt * (1.0 - color.a); - if color.a > FULL_ALPHA { break; } - } - var pos = (pos_view + dir_view * real_t).xyz; - pos[axis] = round(pos[axis]) - (1.0 - dir_uf[axis]); - let vcolor = get_color(node & LEAF_MASK, pos); - var normal = normals[axis]; - let light_color = vec3(1.0); - let light_dir = global_lights[0].dir; - - let diffuse = max(dot(light_dir, normal), 0.0) * light_color; - let ambient = AMBIENT * light_color; - let spec_val = pow(max(dot(dir_view.xyz, reflect(-light_dir, normal)), 0.0), 32.0) * SPECULAR; - let specular = spec_val * light_color; - let new_color = (ambient + diffuse + specular) * vcolor.xyz; - let new_a = min(vcolor.a + spec_val, 1.0); - - color += vec4(new_color.xyz * new_a, new_a) * (1.0 - color.a); - if color.a > FULL_ALPHA { break; } - } + let id = node & LEAF_MASK; + hits[depth] = RayHit(t, id, axis); + min_alpha += min_alpha(id) * (1.0 - min_alpha); + depth += 1u; prev = node; + if depth == 10 || min_alpha >= FULL_ALPHA { break; } } // move to next time point and determine which axis to move along @@ -226,10 +223,15 @@ fn trace_full(pos_view: vec4, dir_view: vec4) -> vec4 { node_start = node; } } - // let fog = min(t / t_mult / 1000.0, 1.0); - // return vec4(color.xyz * (1.0 - fog) + vec3(fog), color.a * (1.0 - fog) + fog); - // return vec4(f32(iters) / f32(MAX_ITERS), 0.0, 0.0, 1.0); - return color; + return RayResult(hits, depth); +} + +fn trace_chunk( + offset: u32, + inv_dir_bits: u32, + t: f32, t_mult: f32, inc_t: vec3, + min_adj: vec3 +) { } fn dir_to_vec(bits: u32) -> vec3 { @@ -240,29 +242,53 @@ fn vec_to_dir(vec: vec3) -> u32 { return vec.x * 4 + vec.y * 2 + vec.z * 1; } -fn get_color(id: u32, pos: vec3) -> vec4 { +fn min_alpha(id: u32) -> f32 { + switch id { + case 0u: {return 0.0;} + case 3u: {return 0.5;} + default: {return 1.0;} + } +} + +const AMBIENT: f32 = 0.2; +const SPECULAR: f32 = 0.5; + +// returns premultiplied +fn shade(id: u32, pos: vec3, normal: vec3, dir_view: vec3, dist: f32) -> vec4 { + var color = vec4(0.0); + if id == 0 { + return color; + } let random = random(floor(pos)); let random2 = random(floor(pos) + vec3(0.0001)); switch id { case 0u: { - return vec4(0.0); + color = vec4(0.0); } case 1u: { - let color = vec3(0.5, 0.5, 0.5 + random * 0.2) * (random2 * 0.4 + 0.8); - return vec4(color, 1.0); + color = vec4(vec3(0.5, 0.5, 0.5 + random * 0.2) * (random2 * 0.4 + 0.8), 1.0); } case 2u: { - let color = vec3(0.4 + random * 0.2, 0.9, 0.4 + random * 0.2) * (random2 * 0.2 + 0.9); - return vec4(color, 1.0); + color = vec4(vec3(0.4 + random * 0.2, 0.9, 0.4 + random * 0.2) * (random2 * 0.2 + 0.9), 1.0); } case 3u: { - let color = vec3(0.5, 0.5, 1.0) * (random2 * 0.2 + 0.8); - return vec4(color, 0.5); - } - default: { - return vec4(1.0, 0.0, 0.0, 1.0); + let fog = min(dist / 64.0, 1.0); + let a = 0.5; + let rgb = vec3(0.5, 0.5, 1.0) * (random2 * 0.2 + 0.8); + color = vec4(rgb * (1.0 - fog * a), a + fog * (1.0 - a)); } + default: {} } + let light_color = vec3(1.0); + let light_dir = global_lights[0].dir; + + let diffuse = max(dot(light_dir, normal), 0.0) * light_color; + let ambient = AMBIENT * light_color; + let spec_val = pow(max(dot(dir_view.xyz, reflect(-light_dir, normal)), 0.0), 32.0) * SPECULAR; + let specular = spec_val * light_color; + let new_color = (ambient + diffuse + specular) * color.xyz; + let new_a = min(color.a + spec_val, 1.0); + return vec4(new_color * new_a, new_a); } fn random(pos: vec3) -> f32 { diff --git a/src/client/render/voxel/ray_oct/shader/compute_shadow.wgsl b/src/client/render/voxel/ray_oct/shader/compute_shadow.wgsl deleted file mode 100644 index 4a60e9f..0000000 --- a/src/client/render/voxel/ray_oct/shader/compute_shadow.wgsl +++ /dev/null @@ -1,419 +0,0 @@ -@group(0) @binding(0) -var view: View; -@group(0) @binding(1) -var voxels: array; -@group(0) @binding(2) -var voxel_groups: array; -@group(0) @binding(3) -var global_lights: array; -@group(0) @binding(4) -var output: texture_storage_2d; - -struct GlobalLight { - dir: vec3, -}; - -struct View { - transform: mat4x4, - zoom: f32, -}; - -struct VoxelGroup { - transform: mat4x4, - transform_inv: mat4x4, - scale: u32, - offset: u32, -}; - -@compute -@workgroup_size(8, 8, 1) -fn main(@builtin(global_invocation_id) cell: vec3) { - let view_dim = textureDimensions(output); - // get position of the pixel; eye at origin, pixel on plane z = 1 - if cell.x >= view_dim.x || cell.y >= view_dim.y { - return; - } - let view_dim_f = vec2(view_dim); - let aspect = view_dim_f.y / view_dim_f.x; - let pixel_pos = vec3( - (vec2(cell.xy) / view_dim_f - vec2(0.5)) * vec2(2.0, -2.0 * aspect), - view.zoom - ); - let pos = view.transform * vec4(pixel_pos, 1.0); - let dir = view.transform * vec4(normalize(pixel_pos), 0.0); - - var color = trace_full(pos, dir); - let light_mult = clamp((-dot(dir.xyz, global_lights[0].dir) - 0.99) * 200.0, 0.0, 1.0); - let sun_color = light_mult * vec3(1.0, 1.0, 1.0); - let sky_bg = vec3(0.3, 0.6, 1.0); - let sky_color = sun_color + sky_bg * (1.0 - light_mult); - color += vec4(sky_color * (1.0 - color.a), 1.0 - color.a); - color.a = 1.0; - textureStore(output, cell.xy, color); -} - -const LEAF_BIT = 1u << 31u; -const LEAF_MASK = ~LEAF_BIT; - -const ZERO3F = vec3(0.0); -const ZERO2F = vec2(0.0); -const FULL_ALPHA = 0.999; -const EPSILON = 0.00000000001; -const MAX_ITERS = 2000; -// NOTE: CANNOT GO HIGHER THAN 23 due to how floating point -// numbers are stored and the bit manipulation used -const MAX_SCALE: u32 = 10; - -fn trace_full(pos_view: vec4, dir_view: vec4) -> vec4 { - let gi = 0; - let group = voxel_groups[gi]; - if group.scale == 0 { - return vec4(0.0); - } - let dimensions = vec3(1u << group.scale); - let dim_f = vec3(dimensions); - let dim_i = vec3(dimensions); - - // transform so that group is at 0,0 - let pos_start = (group.transform_inv * pos_view).xyz; - var dir = (group.transform_inv * dir_view).xyz; - if dir.x == 0 {dir.x = EPSILON;} - if dir.y == 0 {dir.y = EPSILON;} - if dir.z == 0 {dir.z = EPSILON;} - - let dir_if = sign(dir); - let dir_uf = max(dir_if, vec3(0.0)); - - - - // calculate normals - var normals = mat3x3( - (group.transform * vec4(dir_if.x, 0.0, 0.0, 0.0)).xyz, - (group.transform * vec4(0.0, dir_if.y, 0.0, 0.0)).xyz, - (group.transform * vec4(0.0, 0.0, dir_if.z, 0.0)).xyz, - ); - var axis = 0u; - - // find where ray intersects with group - let pos_min = (vec3(1.0) - dir_uf) * dim_f; - // time of intersection; x = td + p, solve for t - var t_min = (pos_min - pos_start) / dir; - if outside3f(pos_start, ZERO3F, dim_f) { - // points of intersection - let px = pos_start + t_min.x * dir; - let py = pos_start + t_min.y * dir; - let pz = pos_start + t_min.z * dir; - - // check if point is in bounds - let hit = vec3( - inside2f(px.yz, ZERO2F, dim_f.yz), - inside2f(py.xz, ZERO2F, dim_f.xz), - inside2f(pz.xy, ZERO2F, dim_f.xy), - ) && (t_min > ZERO3F); - if !any(hit) { - return vec4(0.0); - } - axis = select(select(2u, 1u, hit.y), 0u, hit.x); - } - let t_mult = f32(1u << (MAX_SCALE - group.scale)); - t_min *= t_mult; - // time to move 1 unit in each direction - let full = f32(1u << MAX_SCALE); - let inc_t = abs(1.0 / dir) * full; - let t_offset = max(max(t_min.x, t_min.y), t_min.z); - var t = max(0.0, t_offset); - - let dir_i = vec3(dir_if); - let dir_u = vec3((dir_i + vec3(1)) / 2); - let dir_bits = vec_to_dir(dir_u); - let inv_dir_bits = 7 - dir_bits; - - var node_start = 1u; - var scale = MAX_SCALE - 1; - var scale_exp2 = 0.5; - var skip = LEAF_BIT; - var color = vec4(0.0); - var parents = array(); - - var child = 0u; - var vox_pos = vec3(1.0); - let t_center = t_min + scale_exp2 * inc_t; - if t > t_center.x { vox_pos.x = 1.5; child |= 4u; } - if t > t_center.y { vox_pos.y = 1.5; child |= 2u; } - if t > t_center.z { vox_pos.z = 1.5; child |= 1u; } - let min_adj = t_min - inc_t; - - var iters = 0; - loop { - if iters == MAX_ITERS { - return vec4(1.0, 0.0, 1.0, 1.0); - } - iters += 1; - let t_corner = vox_pos * inc_t + min_adj; - let node = voxels[group.offset + node_start + (child ^ inv_dir_bits)]; - if node >= LEAF_BIT { - if node != skip && node != LEAF_BIT { - skip = node; - let normal = normals[axis]; - let sun_dir = global_lights[0].dir; - let new_pos = pos_view + dir_view * t / t_mult - vec4(normals[axis] * 0.001, 0.0); - - let light = trace_light(new_pos, vec4(-sun_dir, 0.0)); - let diffuse = max(dot(sun_dir, normal) + 0.1, 0.0); - let ambient = 0.2; - let specular = (exp(max( - -(dot(reflect(dir_view.xyz, normal), sun_dir) + 0.90) * 4.0, 0.0 - )) - 1.0) * light; - let lighting = max(diffuse * light.a, ambient); - - let vcolor = get_color(node & LEAF_MASK); - let new_rgb = min(vcolor.xyz * lighting + specular.xyz + light.xyz * vcolor.xyz, vec3(1.0)); - let new_a = min(vcolor.a + specular.a, 1.0); - let new_color = vec4(new_rgb, new_a); - color += vec4(new_color.xyz * new_color.a, new_color.a) * (1.0 - color.a); - if color.a > FULL_ALPHA { break; } - } - - // move to next time point and determine which axis to move along - let t_next = t_corner + scale_exp2 * inc_t; - t = min(min(t_next.x, t_next.y), t_next.z); - axis = select(select(0u, 1u, t == t_next.y), 2u, t == t_next.z); - let move_dir = 4u >> axis; - - // check if need to pop stack - if (child & move_dir) > 0 { - // calculate new scale; first differing bit after adding - let axis_pos = vox_pos[axis]; - // AWARE - let differing = bitcast(axis_pos) ^ bitcast(axis_pos + scale_exp2); - scale = (bitcast(f32(differing)) >> 23) - 127 - (23 - MAX_SCALE); - scale_exp2 = bitcast((scale + 127 - MAX_SCALE) << 23); - if scale >= MAX_SCALE { break; } - - // restore & recalculate parent - let parent_info = parents[scale]; - node_start = parent_info >> 3; - child = parent_info & 7; - let scale_vec = vec3(scale + 23 - MAX_SCALE); - // remove bits lower than current scale - vox_pos = bitcast>((bitcast>(vox_pos) >> scale_vec) << scale_vec); - } - // move to next child and voxel position - child += move_dir; - vox_pos[axis] += scale_exp2; - } else { - // push current node to stack - parents[scale] = (node_start << 3) + child; - scale -= 1u; - - // calculate child node vars - scale_exp2 *= 0.5; - child = 0u; - let t_center = t_corner + scale_exp2 * inc_t; - if t > t_center.x { vox_pos.x += scale_exp2; child |= 4u; } - if t > t_center.y { vox_pos.y += scale_exp2; child |= 2u; } - if t > t_center.z { vox_pos.z += scale_exp2; child |= 1u; } - node_start += 8 + node; - } - } - // return vec4(f32(iters) / f32(MAX_ITERS), 0.0, 0.0, 1.0); - return color; -} - -fn trace_light(pos_view: vec4, dir_view: vec4) -> vec4 { - let gi = 0; - let group = voxel_groups[gi]; - if group.scale == 0 { - return vec4(0.0); - } - let dimensions = vec3(1u << group.scale); - let dim_f = vec3(dimensions); - let dim_i = vec3(dimensions); - - // transform so that group is at 0,0 - let pos_start = (group.transform_inv * pos_view).xyz; - var dir = (group.transform_inv * dir_view).xyz; - if dir.x == 0 {dir.x = EPSILON;} - if dir.y == 0 {dir.y = EPSILON;} - if dir.z == 0 {dir.z = EPSILON;} - - let dir_if = sign(dir); - let dir_uf = max(dir_if, vec3(0.0)); - - - - // calculate normals - var normals = mat3x3( - (group.transform * vec4(dir_if.x, 0.0, 0.0, 0.0)).xyz, - (group.transform * vec4(0.0, dir_if.y, 0.0, 0.0)).xyz, - (group.transform * vec4(0.0, 0.0, dir_if.z, 0.0)).xyz, - ); - var axis = 0u; - - // find where ray intersects with group - let pos_min = (vec3(1.0) - dir_uf) * dim_f; - // time of intersection; x = td + p, solve for t - var t_min = (pos_min - pos_start) / dir; - if outside3f(pos_start, ZERO3F, dim_f) { - // points of intersection - let px = pos_start + t_min.x * dir; - let py = pos_start + t_min.y * dir; - let pz = pos_start + t_min.z * dir; - - // check if point is in bounds - let hit = vec3( - inside2f(px.yz, ZERO2F, dim_f.yz), - inside2f(py.xz, ZERO2F, dim_f.xz), - inside2f(pz.xy, ZERO2F, dim_f.xy), - ) && (t_min > ZERO3F); - if !any(hit) { - return vec4(0.0); - } - axis = select(select(2u, 1u, hit.y), 0u, hit.x); - } - let t_mult = f32(1u << (MAX_SCALE - group.scale)); - t_min *= t_mult; - // time to move 1 unit in each direction - let full = f32(1u << MAX_SCALE); - let inc_t = abs(1.0 / dir) * full; - let t_offset = max(max(t_min.x, t_min.y), t_min.z); - var t = max(0.0, t_offset); - var old_t = t; - - let dir_i = vec3(dir_if); - let dir_u = vec3((dir_i + vec3(1)) / 2); - let dir_bits = vec_to_dir(dir_u); - let inv_dir_bits = 7 - dir_bits; - - var node_start = 1u; - var scale = MAX_SCALE - 1; - var scale_exp2 = 0.5; - var mask = vec4(0.0); - var skip = LEAF_BIT; - var parents = array(); - - var child = 0u; - var vox_pos = vec3(1.0); - let t_center = t_min + scale_exp2 * inc_t; - if t > t_center.x { vox_pos.x = 1.5; child |= 4u; } - if t > t_center.y { vox_pos.y = 1.5; child |= 2u; } - if t > t_center.z { vox_pos.z = 1.5; child |= 1u; } - let min_adj = t_min - inc_t; - var data = 0u; - - var iters = 0; - loop { - if iters == MAX_ITERS { - return vec4(1.0, 0.0, 1.0, 1.0); - } - iters += 1; - let t_corner = vox_pos * inc_t + min_adj; - let node = voxels[group.offset + node_start + (child ^ inv_dir_bits)]; - if node >= LEAF_BIT { - if node != skip && node != LEAF_BIT { - skip = node; - if data == 3 { - let dist = (t - old_t) / t_mult; - let vcolor = vec4(vec3(0.0), min(dist / 12.0, 1.0)); - mask += vec4(vcolor.xyz * vcolor.a, vcolor.a) * (1.0 - mask.a); - } - data = node & LEAF_MASK; - if data != 3 && data != 0 { - let vcolor = get_color(data); - mask += vec4(vcolor.xyz * vcolor.a, vcolor.a) * (1.0 - mask.a); - } - old_t = t; - if mask.a > FULL_ALPHA { break; } - } - - // move to next time point and determine which axis to move along - let t_next = t_corner + scale_exp2 * inc_t; - t = min(min(t_next.x, t_next.y), t_next.z); - axis = select(select(0u, 1u, t == t_next.y), 2u, t == t_next.z); - let move_dir = 4u >> axis; - - // check if need to pop stack - if (child & move_dir) > 0 { - // calculate new scale; first differing bit after adding - let axis_pos = vox_pos[axis]; - // AWARE - let differing = bitcast(axis_pos) ^ bitcast(axis_pos + scale_exp2); - scale = (bitcast(f32(differing)) >> 23) - 127 - (23 - MAX_SCALE); - scale_exp2 = bitcast((scale + 127 - MAX_SCALE) << 23); - if scale >= MAX_SCALE { break; } - - // restore & recalculate parent - let parent_info = parents[scale]; - node_start = parent_info >> 3; - child = parent_info & 7; - let scale_vec = vec3(scale + 23 - MAX_SCALE); - // remove bits lower than current scale - vox_pos = bitcast>((bitcast>(vox_pos) >> scale_vec) << scale_vec); - } - // move to next child and voxel position - child += move_dir; - vox_pos[axis] += scale_exp2; - } else { - // push current node to stack - parents[scale] = (node_start << 3) + child; - scale -= 1u; - - // calculate child node vars - scale_exp2 *= 0.5; - child = 0u; - let t_center = t_corner + scale_exp2 * inc_t; - if t > t_center.x { vox_pos.x += scale_exp2; child |= 4u; } - if t > t_center.y { vox_pos.y += scale_exp2; child |= 2u; } - if t > t_center.z { vox_pos.z += scale_exp2; child |= 1u; } - node_start += 8 + node; - } - } - if data == 3 { - let dist = (t - old_t) / t_mult; - let vcolor = vec4(vec3(0.0), min(dist / 12.0, 1.0)); - mask += vec4(vcolor.xyz * vcolor.a, vcolor.a) * (1.0 - mask.a); - } - mask.a = 1.0 - mask.a; - mask = vec4(mask.a * mask.xyz, mask.a); - return mask; -} - -fn dir_to_vec(bits: u32) -> vec3 { - return vec3(bits >> 2, (bits & 2) >> 1, bits & 1); -} - -fn vec_to_dir(vec: vec3) -> u32 { - return vec.x * 4 + vec.y * 2 + vec.z * 1; -} - -fn get_color(id: u32) -> vec4 { - switch id { - case 0u: { - return vec4(0.0); - } - case 1u: { - return vec4(0.5, 0.5, 0.5, 1.0); - } - case 2u: { - return vec4(0.5, 1.0, 0.5, 1.0); - } - case 3u: { - return vec4(0.5, 0.5, 1.0, 0.5); - } - default: { - return vec4(1.0, 0.0, 0.0, 1.0); - } - } -} - -fn outside3f(v: vec3, low: vec3, high: vec3) -> bool { - return any(v < low) || any(v > high); -} - -fn inside2f(v: vec2, low: vec2, high: vec2) -> bool { - return all(v >= low) && all(v <= high); -} - -fn inside3i(v: vec3, low: vec3, high: vec3) -> bool { - return all(v >= low) && all(v <= high); -} diff --git a/src/client/render/voxel/ray_oct/shader/compute_struct.wgsl b/src/client/render/voxel/ray_oct/shader/compute_single.wgsl similarity index 54% rename from src/client/render/voxel/ray_oct/shader/compute_struct.wgsl rename to src/client/render/voxel/ray_oct/shader/compute_single.wgsl index 3554f06..91e2890 100644 --- a/src/client/render/voxel/ray_oct/shader/compute_struct.wgsl +++ b/src/client/render/voxel/ray_oct/shader/compute_single.wgsl @@ -35,67 +35,13 @@ fn main(@builtin(global_invocation_id) cell: vec3) { } let view_dim_f = vec2(view_dim); let aspect = view_dim_f.y / view_dim_f.x; - let pixel_pos = vec3( - (vec2(cell.xy) / view_dim_f - vec2(0.5)) * vec2(2.0, -2.0 * aspect), - view.zoom + let pixel_pos = vec2( + (vec2(cell.xy) / view_dim_f - vec2(0.5)) * vec2(2.0, -2.0 * aspect) ); - let pos = view.transform * vec4(pixel_pos, 1.0); - let dir = view.transform * vec4(normalize(pixel_pos), 0.0); + let pos = view.transform * vec4(pixel_pos, 1.0, 1.0); + let dir = view.transform * vec4(normalize(vec3(pixel_pos, view.zoom)), 0.0); - let start = start_ray(pos, dir); - var color = vec4(0.0); - let ambient = 0.2; - - if start.hit { - var res = ray_next(start.ray, LEAF_BIT); - var normals = start.normals; - let specular = (exp(max( - -(dot(reflect(dir.xyz, normals[res.ray.axis]), global_lights[0].dir) + 0.90) * 4.0, 0.0 - )) - 1.0); - while res.data != 0 { - let data = res.data & LEAF_MASK; - if data != 0 { - let vcolor = get_color(data); - let diffuse = max(dot(global_lights[0].dir, normals[res.ray.axis]) + 0.1, 0.0); - let light = max(diffuse, ambient); - let new_color = min(vcolor.xyz * light, vec3(1.0)); - color += vec4(new_color.xyz * vcolor.a, vcolor.a) * (1.0 - color.a); - if color.a > FULL_ALPHA { break; } - } - let old_t = res.ray.t; - res = ray_next(res.ray, res.data); - if data == 3 { - let dist = (res.ray.t - old_t) / start.t_mult; - let a = min(dist / 12.0, 1.0); - color += vec4(vec3(0.0) * a, a) * (1.0 - color.a); - } - } - if color.a != 0 { - let pos = pos + dir * res.ray.t / start.t_mult - vec4(normals[res.ray.axis] * 0.001, 0.0); - let dir = vec4(-global_lights[0].dir, 0.0); - let start = start_ray(pos, dir); - res = ray_next(start.ray, LEAF_BIT); - var light = 1.0; - while res.data != 0 { - let data = res.data & LEAF_MASK; - if data != 0 { - let vcolor = get_color(data); - if data != 3 { light -= vcolor.a * light; } - if light <= 0 { break; } - } - let old_t = res.ray.t; - res = ray_next(res.ray, res.data); - if data == 3 { - let dist = (res.ray.t - old_t) / start.t_mult; - let a = min(dist / 12.0, 1.0); - light -= a; - } - } - color = vec4(color.xyz * max(light, ambient), color.a) + vec4(vec3(specular * light), 0.0); - } - // color = vec4(pos.xyz / 128.0, 1.0); - } - // var color = trace_full(pos, dir); + var color = trace_full(pos, dir); let light_mult = clamp((-dot(dir.xyz, global_lights[0].dir) - 0.99) * 200.0, 0.0, 1.0); let sun_color = light_mult * vec3(1.0, 1.0, 1.0); let sky_bg = vec3(0.3, 0.6, 1.0); @@ -112,49 +58,25 @@ const ZERO3F = vec3(0.0); const ZERO2F = vec2(0.0); const FULL_ALPHA = 0.999; const EPSILON = 0.00000000001; -const MAX_ITERS = 2000; +const MAX_ITERS = 10000; // NOTE: CANNOT GO HIGHER THAN 23 due to how floating point // numbers are stored and the bit manipulation used -const MAX_SCALE: u32 = 10; +const MAX_SCALE: u32 = 13; +const AMBIENT: f32 = 0.2; +const SPECULAR: f32 = 0.5; -struct Ray { - t: f32, - vox_pos: vec3, - t_inc: vec3, - scale: u32, - min_adj: vec3, - child: u32, - axis: u32, - node_start: u32, - group_offset: u32, - inv_dir_bits: u32, - parents: array, -}; - -struct RayResult { - ray: Ray, - data: u32, -} - -struct RayStart { - hit: bool, - ray: Ray, - normals: mat3x3, - t_mult: f32, -} - -fn start_ray(pos_view: vec4, dir_view: vec4) -> RayStart { +fn trace_full(pos_view: vec4, dir_view: vec4) -> vec4 { let gi = 0; let group = voxel_groups[gi]; if group.scale == 0 { - return RayStart(); + return vec4(0.0); } let dimensions = vec3(1u << group.scale); let dim_f = vec3(dimensions); let dim_i = vec3(dimensions); // transform so that group is at 0,0 - let pos = (group.transform_inv * pos_view).xyz; + let pos_start = (group.transform_inv * pos_view).xyz; var dir = (group.transform_inv * dir_view).xyz; if dir.x == 0 {dir.x = EPSILON;} if dir.y == 0 {dir.y = EPSILON;} @@ -176,12 +98,12 @@ fn start_ray(pos_view: vec4, dir_view: vec4) -> RayStart { // find where ray intersects with group let pos_min = (vec3(1.0) - dir_uf) * dim_f; // time of intersection; x = td + p, solve for t - var t_min = (pos_min - pos) / dir; - if outside3f(pos, ZERO3F, dim_f) { + var t_min = (pos_min - pos_start) / dir; + if outside3f(pos_start, ZERO3F, dim_f) { // points of intersection - let px = pos + t_min.x * dir; - let py = pos + t_min.y * dir; - let pz = pos + t_min.z * dir; + let px = pos_start + t_min.x * dir; + let py = pos_start + t_min.y * dir; + let pz = pos_start + t_min.z * dir; // check if point is in bounds let hit = vec3( @@ -190,81 +112,80 @@ fn start_ray(pos_view: vec4, dir_view: vec4) -> RayStart { inside2f(pz.xy, ZERO2F, dim_f.xy), ) && (t_min > ZERO3F); if !any(hit) { - return RayStart(); + return vec4(0.0); } axis = select(select(2u, 1u, hit.y), 0u, hit.x); } let t_mult = f32(1u << (MAX_SCALE - group.scale)); t_min *= t_mult; - // time to move 1 unit in each direction let full = f32(1u << MAX_SCALE); - let t_inc = abs(1.0 / dir) * full; + // time to move entire side length in each direction + let inc_t = abs(1.0 / dir) * full; let t_offset = max(max(t_min.x, t_min.y), t_min.z); - let t = max(0.0, t_offset); + var t = max(0.0, t_offset); let dir_i = vec3(dir_if); - let dir_u = vec3((dir_i + vec3(1)) / 2); + let dir_u = vec3(dir_uf); let dir_bits = vec_to_dir(dir_u); let inv_dir_bits = 7 - dir_bits; - let node_start = 1u; - let scale = MAX_SCALE - 1; - let scale_exp2 = 0.5; - let parents = array(); + var node_start = 1u; + var scale = MAX_SCALE - 1; + var scale_exp2 = 0.5; + var color = vec4(0.0); + var parents = array(); + var prev = LEAF_BIT; + var old_t = t / t_mult; var child = 0u; var vox_pos = vec3(1.0); - let t_center = t_min + scale_exp2 * t_inc; + let t_center = t_min + scale_exp2 * inc_t; if t > t_center.x { vox_pos.x = 1.5; child |= 4u; } if t > t_center.y { vox_pos.y = 1.5; child |= 2u; } if t > t_center.z { vox_pos.z = 1.5; child |= 1u; } - let min_adj = t_min - t_inc; + let min_adj = t_min - inc_t; - return RayStart( - true, - Ray( - t, - vox_pos, - t_inc, - scale, - min_adj, - child, - axis, - node_start, - group.offset, - inv_dir_bits, - parents, - ), - normals, - t_mult, - ); -} - -fn ray_next(ray: Ray, skip: u32) -> RayResult { - let group_offset = ray.group_offset; - let t_inc = ray.t_inc; - let min_adj = ray.min_adj; - let inv_dir_bits = ray.inv_dir_bits; - var scale = ray.scale; - var scale_exp2 = bitcast((scale + 127 - MAX_SCALE) << 23); - var vox_pos = ray.vox_pos; - var t = ray.t; - var node_start = ray.node_start; - var child = ray.child; - var parents = ray.parents; - var axis: u32; - var data = 0u; + var iters = 0; loop { - let t_corner = vox_pos * t_inc + min_adj; - let node = voxels[group_offset + node_start + (child ^ inv_dir_bits)]; + if iters == MAX_ITERS { + return vec4(1.0, 0.0, 1.0, 1.0); + } + iters += 1; + let t_corner = vox_pos * inc_t + min_adj; + let node = voxels[group.offset + node_start + (child ^ inv_dir_bits)]; if node >= LEAF_BIT { - if node != skip { - data = node; - break; + if node != prev { + if node != LEAF_BIT { + let real_t = t / t_mult; + let dist = real_t - old_t; + old_t = real_t; + let filt = min(dist / 64.0, 1.0); + if prev == LEAF_BIT + 3 { + color.a += filt * (1.0 - color.a); + if color.a > FULL_ALPHA { break; } + } + var pos = (pos_view + dir_view * real_t).xyz; + pos[axis] = round(pos[axis]) - (1.0 - dir_uf[axis]); + let vcolor = get_color(node & LEAF_MASK, pos); + var normal = normals[axis]; + let light_color = vec3(1.0); + let light_dir = global_lights[0].dir; + + let diffuse = max(dot(light_dir, normal), 0.0) * light_color; + let ambient = AMBIENT * light_color; + let spec_val = pow(max(dot(dir_view.xyz, reflect(-light_dir, normal)), 0.0), 32.0) * SPECULAR; + let specular = spec_val * light_color; + let new_color = (ambient + diffuse + specular) * vcolor.xyz; + let new_a = min(vcolor.a + spec_val, 1.0); + + color += vec4(new_color.xyz * new_a, new_a) * (1.0 - color.a); + if color.a > FULL_ALPHA { break; } + } + prev = node; } // move to next time point and determine which axis to move along - let t_next = t_corner + scale_exp2 * t_inc; + let t_next = t_corner + scale_exp2 * inc_t; t = min(min(t_next.x, t_next.y), t_next.z); axis = select(select(0u, 1u, t == t_next.y), 2u, t == t_next.z); let move_dir = 4u >> axis; @@ -298,29 +219,17 @@ fn ray_next(ray: Ray, skip: u32) -> RayResult { // calculate child node vars scale_exp2 *= 0.5; child = 0u; - let t_center = t_corner + scale_exp2 * t_inc; + let t_center = t_corner + scale_exp2 * inc_t; if t > t_center.x { vox_pos.x += scale_exp2; child |= 4u; } if t > t_center.y { vox_pos.y += scale_exp2; child |= 2u; } if t > t_center.z { vox_pos.z += scale_exp2; child |= 1u; } - node_start += 8 + node; + node_start = node; } } - return RayResult( - Ray( - t, - vox_pos, - t_inc, - scale, - min_adj, - child, - axis, - node_start, - group_offset, - inv_dir_bits, - parents, - ), - data - ); + // let fog = min(t / t_mult / 1000.0, 1.0); + // return vec4(color.xyz * (1.0 - fog) + vec3(fog), color.a * (1.0 - fog) + fog); + // return vec4(f32(iters) / f32(MAX_ITERS), 0.0, 0.0, 1.0); + return color; } fn dir_to_vec(bits: u32) -> vec3 { @@ -331,19 +240,24 @@ fn vec_to_dir(vec: vec3) -> u32 { return vec.x * 4 + vec.y * 2 + vec.z * 1; } -fn get_color(id: u32) -> vec4 { +fn get_color(id: u32, pos: vec3) -> vec4 { + let random = random(floor(pos)); + let random2 = random(floor(pos) + vec3(0.0001)); switch id { case 0u: { return vec4(0.0); } case 1u: { - return vec4(0.5, 0.5, 0.5, 1.0); + let color = vec3(0.5, 0.5, 0.5 + random * 0.2) * (random2 * 0.4 + 0.8); + return vec4(color, 1.0); } case 2u: { - return vec4(0.5, 1.0, 0.5, 1.0); + let color = vec3(0.4 + random * 0.2, 0.9, 0.4 + random * 0.2) * (random2 * 0.2 + 0.9); + return vec4(color, 1.0); } case 3u: { - return vec4(0.5, 0.5, 1.0, 0.5); + let color = vec3(0.5, 0.5, 1.0) * (random2 * 0.2 + 0.8); + return vec4(color, 0.5); } default: { return vec4(1.0, 0.0, 0.0, 1.0); @@ -351,6 +265,10 @@ fn get_color(id: u32) -> vec4 { } } +fn random(pos: vec3) -> f32 { + return fract(sin(dot(pos, vec3(12.9898, 78.233, 25.1279))) * 43758.5453123); +} + fn outside3f(v: vec3, low: vec3, high: vec3) -> bool { return any(v < low) || any(v > high); } diff --git a/src/client/render/voxel/ray_oct/view.rs b/src/client/render/voxel/ray_oct/view.rs index 6419a7d..d07f65a 100644 --- a/src/client/render/voxel/ray_oct/view.rs +++ b/src/client/render/voxel/ray_oct/view.rs @@ -1,10 +1,14 @@ use nalgebra::Transform3; +use crate::common::component::chunk::SCALE; + #[repr(C, align(16))] #[derive(Clone, Copy, PartialEq, bytemuck::Zeroable)] pub struct View { pub transform: Transform3, pub zoom: f32, + pub chunk_scale: u32, + pub chunk_radius: u32, } unsafe impl bytemuck::Pod for View {} @@ -14,6 +18,8 @@ impl Default for View { Self { zoom: 1.0, transform: Transform3::identity(), + chunk_scale: SCALE, + chunk_radius: 2, } } } diff --git a/src/common/component/chunk/mod.rs b/src/common/component/chunk/mod.rs index 6946910..10a64c2 100644 --- a/src/common/component/chunk/mod.rs +++ b/src/common/component/chunk/mod.rs @@ -8,7 +8,7 @@ use bevy_derive::{Deref, DerefMut}; use bevy_ecs::{bundle::Bundle, component::Component, entity::Entity, system::Resource}; use nalgebra::Vector3; -pub const SCALE: u32 = 13; +pub const SCALE: u32 = 9; pub const SIDE_LENGTH: usize = 2usize.pow(SCALE); pub const SHAPE: (usize, usize, usize) = (SIDE_LENGTH, SIDE_LENGTH, SIDE_LENGTH); pub const DIMENSIONS: Vector3 = Vector3::new(SIDE_LENGTH, SIDE_LENGTH, SIDE_LENGTH); diff --git a/src/util/oct_tree.rs b/src/util/oct_tree.rs index 5935523..fa57c9b 100644 --- a/src/util/oct_tree.rs +++ b/src/util/oct_tree.rs @@ -86,7 +86,7 @@ impl OctTree { assert!(levels > 0); let mut data = Vec::new(); let mut map = OctNodeMap::default(); - data.push(OctNode::new_node(0)); + data.push(OctNode::new_node(1)); Self::from_fn_offset_inner(f_leaf, f_node, &mut data, levels, offset, &mut map); if data.len() == 2 { data.remove(0); diff --git a/src/util/timer.rs b/src/util/timer.rs index 38635d6..1d8d0e3 100644 --- a/src/util/timer.rs +++ b/src/util/timer.rs @@ -25,6 +25,11 @@ impl Timer { self.times[self.pos] = Some(self.start); self.pos = (self.pos + 1) % self.times.len(); } + pub fn add(&mut self, duration: Duration) { + self.durs[self.pos] = Some(duration); + self.times[self.pos] = Some(Instant::now()); + self.pos = (self.pos + 1) % self.times.len(); + } pub fn avg(&self) -> Duration { let filtered: Vec<_> = self.durs.iter().filter_map(|d| *d).collect(); let len = filtered.len(); @@ -45,8 +50,9 @@ impl Timer { pub fn per_sec(&self) -> usize { let now = Instant::now(); let mut count = 0; - while count < self.times.len() { - let i = (self.pos + count + 1) % self.times.len(); + let len = self.times.len(); + while count < len { + let i = (self.pos + len - count - 1) % len; let Some(t) = self.times[i] else { break }; if now - t <= Duration::from_secs(1) { count += 1;