From f8cfb69b216eee205fe78be0e111217db8632842 Mon Sep 17 00:00:00 2001 From: Albin Chaboissier Date: Sat, 5 Sep 2026 15:39:19 +0200 Subject: [PATCH] Works --- Cargo.toml | 1 + shaders/voxel.wgsl | 13 +- src/main.rs | 88 ++- src/producers.rs | 191 +++++- src/voxel/pipeline.rs | 799 ----------------------- src/{voxel.rs => voxel_cache.rs} | 0 src/{voxel => voxel_cache}/cache.rs | 289 +++++--- src/{voxel => voxel_cache}/cache.wait.rs | 0 src/{voxel => voxel_cache}/gpu.rs | 2 +- src/voxel_cache/pipeline.rs | 68 ++ src/voxel_cache/request_buffer.rs | 0 src/{voxel => voxel_cache}/sparse.rs | 4 +- src/voxel_cache/structure_table.rs | 0 src/voxel_cache/usage_buffer.rs | 0 14 files changed, 495 insertions(+), 960 deletions(-) delete mode 100644 src/voxel/pipeline.rs rename src/{voxel.rs => voxel_cache.rs} (100%) rename src/{voxel => voxel_cache}/cache.rs (89%) rename src/{voxel => voxel_cache}/cache.wait.rs (100%) rename src/{voxel => voxel_cache}/gpu.rs (93%) create mode 100644 src/voxel_cache/pipeline.rs create mode 100644 src/voxel_cache/request_buffer.rs rename src/{voxel => voxel_cache}/sparse.rs (99%) create mode 100644 src/voxel_cache/structure_table.rs create mode 100644 src/voxel_cache/usage_buffer.rs diff --git a/Cargo.toml b/Cargo.toml index 382199a..cf25611 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -14,6 +14,7 @@ env_logger = "0.11.11" fastapprox = "0.3.1" glam = "0.33.5" image = "0.25.10" +indicatif = "0.18.6" itertools = "0.15.0" pollster = "1.0.1" rand = "0.10.2" diff --git a/shaders/voxel.wgsl b/shaders/voxel.wgsl index e9a4568..5381edc 100644 --- a/shaders/voxel.wgsl +++ b/shaders/voxel.wgsl @@ -216,9 +216,9 @@ fn new_traverse(ray_dir: vec3, ray_origin: vec3, root_id: u32, dist_of { let max_depth = 5; let dist_offset_voxel = dist_offset * f32(1 << u32(max_depth * 2)); - let fovy_deg = 100.; - let fovy_rad = (fovy_deg / 180.) * 3.14159; - let cone_factor = tan(fovy_rad) * 2.; + let fovy_deg = 100. / 1920.; + let fovy_rad = (fovy_deg * 3.14) / 180.; + let cone_factor = tan(fovy_rad / 2.) * 2.; let st_pointer = structure_table_pointer[root_id]; @@ -272,9 +272,9 @@ fn new_traverse(ray_dir: vec3, ray_origin: vec3, root_id: u32, dist_of var child_pos = (voxel >> vec3(u32(node_shift - 2))) & vec3(3); var pointer = structure_pool[dfs_stack[current_depth]].pointers[child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4]; - let min_child_size = (length(vec3(voxel) - pos_origin) + dist_offset) * cone_factor; + let min_child_size = (length(vec3(voxel) - pos_origin) + dist_offset_voxel) * cone_factor; while(node_subdivided(pointer) && - f32(child_size / 8) >= min_child_size + f32(child_size / 4) >= min_child_size ) { if(!node_pointer_valid(pointer) && node_subdivided(pointer)) @@ -558,10 +558,11 @@ fn fragment(in: VertexOutput) -> FragmentOutput let ray_origin = (in.cam_pos - in.chunk_position) + ray_dir * (max(0., interp.x)); - let result = new_traverse(ray_dir, ray_origin, in.structure_id, length(in.cam_pos - ray_origin)); + let result = new_traverse(ray_dir, ray_origin, in.structure_id, length(in.cam_pos - (ray_origin + in.chunk_position))); let clip_pos = constants.view_proj * vec4(result.hit_pos + in.chunk_position, 1.); let depth = clip_pos.z / clip_pos.w; var frag_out: FragmentOutput; + //frag_out.color = result.color; frag_out.color = result.color; frag_out.depth = depth; return frag_out; diff --git a/src/main.rs b/src/main.rs index a56b7dc..3099845 100644 --- a/src/main.rs +++ b/src/main.rs @@ -1,4 +1,5 @@ #![feature(generic_const_exprs)] +#![feature(float_algebraic)] use core::sync; use std::cell::RefCell; @@ -12,6 +13,7 @@ use std::sync::mpsc::sync_channel; use bytemuck::Pod; use bytemuck::Zeroable; +use bytemuck::cast_slice; use crevice::std140::AsStd140; use crevice::std430::AsStd430; use egui::Color32; @@ -78,26 +80,18 @@ use crate::producers::ChunkedProducer; use crate::producers::Producer; use crate::producers::SineGenerator; use crate::producers::TerrainGenerator; -use crate::voxel::cache::CacheNodeRequest; -use crate::voxel::cache::CacheResponse; -use crate::voxel::cache::ColorBytes; -use crate::voxel::cache::DestinationElement; -use crate::voxel::cache::LocationPoolElement; -use crate::voxel::cache::RequestBuffer; -use crate::voxel::cache::UsageBuffer; -use crate::voxel::cache::VoxelCache; -use crate::voxel::gpu::ExplicitNTreeNode; -use crate::voxel::pipeline::ChunkHandle; -use crate::voxel::pipeline::ChunkObject; -use crate::voxel::pipeline::VoxelPipeline; -use crate::voxel::sparse::Color; -use crate::voxel::sparse::NTree; -use crate::voxel::sparse::NTreeNodeLocator; +use crate::voxel_cache::cache::CacheNodeRequest; +use crate::voxel_cache::cache::CacheResponse; +use crate::voxel_cache::cache::ColorBytes; +use crate::voxel_cache::cache::DestinationElement; +use crate::voxel_cache::cache::LocationPoolElement; +use crate::voxel_cache::cache::VoxelCache; +use crate::voxel_cache::sparse::NTreeNodeLocator; mod camera; mod egui_renderer; mod producers; -mod voxel; +mod voxel_cache; //mod tree; // @@ -120,6 +114,7 @@ struct State instance_buffer: Buffer, instance_count: usize, usage_vec: Arc>>, + rm_time: Arc>, insertion_debounce: bool, camera: Camera, @@ -188,9 +183,21 @@ impl State let egui_renderer = EguiRenderer::new(&device, surface_format, &window); - let mut voxel_cache = VoxelCache::<4>::new(80_000, device.clone(), queue.clone()); + let mut voxel_cache = VoxelCache::<4>::new(200_000, device.clone(), queue.clone()); - let terrain_generator = TerrainGenerator::<4>::new(5, "vxls_height.tif", 0.2, "img.jpg"); + //let terrain_generator = TerrainGenerator::<4>::new(5, "vxls_height.tif", 0.2, "img.jpg"); + // let terrain_generator = TerrainGenerator::<4>::new( + // 5, + // "./pointe_percee/height.tif", + // 0.2, + // "./pointe_percee/ortho.jpg", + // ); + let terrain_generator = TerrainGenerator::<4>::new( + 5, + "/home/albin/Documents/vxls_maps/lapiz/height.tif", + 0.2, + "/home/albin/Documents/vxls_maps/lapiz/ortho.jpg", + ); let mut chunk_pos_map = HashMap::new(); let chunk_instances = (0..terrain_generator.chunk_width) @@ -307,6 +314,7 @@ impl State instance_count, terrain_generator: Arc::new(terrain_generator), chunk_pos_map: chunk_pos_map.into(), + rm_time: Arc::new(Mutex::new(0.)), }; // Configure surface for the first time @@ -427,6 +435,19 @@ impl State } }; + let timestamp_query = self.device.create_query_set(&wgpu::QuerySetDescriptor { + label: Some("timestamp_query_set"), + ty: wgpu::QueryType::Timestamp, + count: 2, + }); + + let timestamp_buffer = self.device.create_buffer(&wgpu::BufferDescriptor { + label: Some("timestamp_buffer"), + size: (size_of::() * 2) as u64, + usage: BufferUsages::QUERY_RESOLVE | BufferUsages::COPY_SRC, + mapped_at_creation: false, + }); + let texture_view = surface_texture .texture .create_view(&wgpu::TextureViewDescriptor { @@ -459,7 +480,11 @@ impl State }), stencil_ops: None, }), - timestamp_writes: None, + timestamp_writes: Some(wgpu::RenderPassTimestampWrites { + query_set: ×tamp_query, + beginning_of_pass_write_index: Some(0), + end_of_pass_write_index: Some(1), + }), occlusion_query_set: None, multiview_mask: None, }); @@ -477,11 +502,13 @@ impl State // End the renderpass. drop(renderpass); + + encoder.resolve_query_set(×tamp_query, 0..2, ×tamp_buffer, 0); } let requests = self.device.create_buffer(&wgpu::BufferDescriptor { label: Some("dummy_dumb_dinky_aaaahhh_buffer"), - size: 16 * 1024, + size: 16 * 2000, usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC, mapped_at_creation: false, }); @@ -513,6 +540,11 @@ impl State .size(28.), ); } + + ui.label(format!( + "Ray-marching time: {}", + *self.rm_time.lock() / 1_000_000. + )); egui_plot::Plot::new("Plot").show(ui, |plot_ui| { plot_ui.bar_chart(BarChart::new( "histo", @@ -582,6 +614,22 @@ impl State // .contains(&winit::keyboard::KeyCode::KeyF)) // && !self.insertion_debounce // { + + // Report frame time + let cloned_rm_time = self.rm_time.clone(); + let cloned_queue = self.queue.clone(); + DownloadBuffer::read_buffer( + &self.device, + &self.queue, + ×tamp_buffer.slice(..), + move |buffer| { + let buffer_slice = buffer.unwrap(); + let slice: &[u64] = cast_slice(&buffer_slice); + let time = (slice[1] - slice[0]) as f32 * cloned_queue.get_timestamp_period(); + *cloned_rm_time.lock() = time; + }, + ); + if !self .camera .pressed_keyset diff --git a/src/producers.rs b/src/producers.rs index 209c85b..aec64f2 100644 --- a/src/producers.rs +++ b/src/producers.rs @@ -2,11 +2,12 @@ use std::fs::File; use std::path::Path; use glam::Vec3; +use indicatif::ProgressIterator; use itertools::Itertools; -use crate::voxel::gpu::ExplicitNTreeNode; -use crate::voxel::gpu::StructurePointer; -use crate::voxel::sparse::Color; +use crate::voxel_cache::gpu::ExplicitNTreeNode; +use crate::voxel_cache::gpu::StructurePointer; +use crate::voxel_cache::sparse::Color; pub struct BallGenerator { @@ -15,7 +16,13 @@ pub struct BallGenerator pub fn map(x: f32, x_min: f32, x_max: f32, y_min: f32, y_max: f32) -> f32 { - ((x - x_min) / (x_max - x_min)) * (y_max - y_min) + y_min + //((x - x_min) / (x_max - x_min)) * (y_max - y_min) + y_min + let input_range = x_max.algebraic_sub(x_min); + let output_range = y_max.algebraic_sub(y_min); + + (x.algebraic_sub(x_min).algebraic_div(input_range)) + .algebraic_mul(output_range) + .algebraic_add(y_min) } pub trait Producer @@ -240,6 +247,14 @@ where heightmap: Vec, colormap: Vec, + heightmap_low_width: usize, + heightmap_low_height: usize, + heightmap_low: Vec<(f32, f32)>, + + colormap_low_width: usize, + colormap_low_height: usize, + colormap_mip: Vec, + pub chunk_width: usize, pub chunk_height: usize, pub chunk_alt: usize, @@ -256,29 +271,50 @@ where color_path: P, ) -> Self { + println!("Starting terrain producer"); + println!("Loading height map."); let mut tiff_dec = tiff::decoder::Decoder::new(File::open(height_path).unwrap()).unwrap(); let (heightmap_width, heightmap_height) = tiff_dec.dimensions().unwrap(); let (heightmap_width, heightmap_height) = (heightmap_width as usize, heightmap_height as usize); - let heightmap = match tiff_dec.read_image().unwrap() + let mut heightmap = match tiff_dec.read_image().unwrap() { tiff::decoder::DecodingResult::F32(vec) => vec, _ => panic!("Unsupported format"), }; + println!("Loading color map."); let mut color = image::ImageReader::open(color_path).unwrap(); color.no_limits(); let color = color.decode().unwrap(); - let colormap = color.as_rgb8().unwrap().to_vec(); + let mut colormap = color.as_rgb8().unwrap().to_vec(); + + println!("Converting color spaces"); + colormap.iter_mut().for_each(|x| { + let normalized = map(*x as f32, 0., 255., 0., 1.); + let maped = normalized.powf(2.4); + *x = map(maped, 0., 1., 0., 255.) as u8; + }); let terrain_width = color.width() as usize; let terrain_height = color.height() as usize; - let heightmap_min = heightmap.iter().copied().reduce(f32::min).unwrap(); + println!("Computing heightmap min/max"); + let heightmap_min = heightmap + .iter() + .copied() + .filter(|x| *x != -9999.) + .reduce(f32::min) + .unwrap(); let heightmap_max = heightmap.iter().copied().reduce(f32::max).unwrap(); + heightmap + .iter_mut() + .filter(|x| **x == -9999.) + .for_each(|x| *x = heightmap_min); + // Decide size in chunks let height_amplitude = heightmap_max - heightmap_min; let chunk_size = N.pow(chunk_power as u32); @@ -286,6 +322,65 @@ where let chunk_height = terrain_height.div_ceil(chunk_size); let chunk_alt = ((height_amplitude / height_factor) as usize).div_ceil(chunk_size); + // build the low heightmap + println!("Computing low res height/color maps"); + let heightmap_low_width = heightmap_width / 8; + let heightmap_low_height = heightmap_height / 8; + + let mut heightmap_low = vec![(0., 0.); heightmap_low_height * heightmap_low_width]; + + for y in 0..heightmap_low_height + { + for x in 0..heightmap_low_width + { + let mut min = heightmap_max; + let mut max = heightmap_min; + for sy in (y * 8)..(y * 8 + 8) + { + for sx in (x * 8)..(x * 8 + 8) + { + min = min.min(heightmap[sx + sy * heightmap_width]); + max = max.max(heightmap[sx + sy * heightmap_width]); + } + } + + heightmap_low[x + y * heightmap_low_width] = (min, max); + } + } + + // build the color map mip + let colormap_low_width = terrain_width / 8; + let colormap_low_height = terrain_height / 8; + + let mut colormap_mip = vec![0u8; colormap_low_width * colormap_low_height * 3]; + + for y in 0..colormap_low_height + { + for x in 0..colormap_low_width + { + let mut r = 0u32; + let mut g = 0u32; + let mut b = 0u32; + for sy in (y * 8)..(y * 8 + 8) + { + for sx in (x * 8)..(x * 8 + 8) + { + r += colormap[(sx + sy * terrain_width) * 3] as u32; + g += colormap[(sx + sy * terrain_width) * 3 + 1] as u32; + b += colormap[(sx + sy * terrain_width) * 3 + 2] as u32; + } + } + + colormap_mip[(x + y * colormap_low_width) * 3] = + ((r as f32) / (8 * 8) as f32).clamp(0., 255.) as u8; + colormap_mip[(x + y * colormap_low_width) * 3 + 1] = + ((g as f32) / (8 * 8) as f32).clamp(0., 255.) as u8; + colormap_mip[(x + y * colormap_low_width) * 3 + 2] = + ((b as f32) / (8 * 8) as f32).clamp(0., 255.) as u8; + } + } + + println!("Producer ready"); Self { chunk_power, heightmap_width, @@ -297,6 +392,14 @@ where heightmap, colormap, + heightmap_low_width, + heightmap_low_height, + heightmap_low, + + colormap_low_width, + colormap_low_height, + colormap_mip, + chunk_width, chunk_height, chunk_alt, @@ -340,30 +443,68 @@ where let mut sample_min = self.heightmap_max; let mut color_avg = Color(0., 0., 0., 0.); let mut count = 0; - for (x, z) in (0..child_size).cartesian_product(0..child_size) + + if depth <= 2 { - let gvx = gcx + x; - let gvz = gcz + z; - if gvx < self.terrain_width && gvz < self.terrain_height + for (z, x) in (0..(child_size / 8)).cartesian_product(0..(child_size / 8)) { - // Height sample - let height_x = (gvx * self.heightmap_width) / self.terrain_width; - let height_z = (gvz * self.heightmap_height) / self.terrain_height; + let gvx = (gcx + x * 8) / 8; + let gvz = (gcz + z * 8) / 8; + if gvx < self.colormap_low_width && gvz < self.colormap_low_height + { + // Height sample + let height_x = (gvx * self.heightmap_low_width) / self.colormap_low_width; + let height_z = (gvz * self.heightmap_low_height) / self.colormap_low_height; - let sample = self.heightmap[height_x + height_z * self.heightmap_width]; - sample_max = sample_max.max(sample); - sample_min = sample_min.min(sample); + let sample = + self.heightmap_low[height_x + height_z * self.heightmap_low_width]; + sample_min = sample_min.min(sample.0); + sample_max = sample_max.max(sample.1); - let sample_color_r = self.colormap[(gvx + gvz * self.terrain_width) * 3]; - let sample_color_g = self.colormap[(gvx + gvz * self.terrain_width) * 3 + 1]; - let sample_color_b = self.colormap[(gvx + gvz * self.terrain_width) * 3 + 2]; + let sample_color_r = + self.colormap_mip[(gvx + gvz * self.colormap_low_width) * 3]; + let sample_color_g = + self.colormap_mip[(gvx + gvz * self.colormap_low_width) * 3 + 1]; + let sample_color_b = + self.colormap_mip[(gvx + gvz * self.colormap_low_width) * 3 + 2]; - color_avg.0 += map(sample_color_r as f32, 0., 256., 0., 1.); - color_avg.1 += map(sample_color_g as f32, 0., 256., 0., 1.); - color_avg.2 += map(sample_color_b as f32, 0., 256., 0., 1.); - count += 1; + color_avg.0 += map(sample_color_r as f32, 0., 256., 0., 1.); + color_avg.1 += map(sample_color_g as f32, 0., 256., 0., 1.); + color_avg.2 += map(sample_color_b as f32, 0., 256., 0., 1.); + count += 1; + } + //let gvy = gcy; + } + } + else + { + for (z, x) in (0..child_size).cartesian_product(0..child_size) + { + let gvx = gcx + x; + let gvz = gcz + z; + if gvx < self.terrain_width && gvz < self.terrain_height + { + // Height sample + let height_x = (gvx * self.heightmap_width) / self.terrain_width; + let height_z = (gvz * self.heightmap_height) / self.terrain_height; + + let sample = self.heightmap[height_x + height_z * self.heightmap_width]; + sample_max = sample_max.max(sample); + sample_min = sample_min.min(sample); + + let sample_color_r = self.colormap[(gvx + gvz * self.terrain_width) * 3]; + let sample_color_g = + self.colormap[(gvx + gvz * self.terrain_width) * 3 + 1]; + let sample_color_b = + self.colormap[(gvx + gvz * self.terrain_width) * 3 + 2]; + + color_avg.0 += map(sample_color_r as f32, 0., 256., 0., 1.); + color_avg.1 += map(sample_color_g as f32, 0., 256., 0., 1.); + color_avg.2 += map(sample_color_b as f32, 0., 256., 0., 1.); + count += 1; + } + //let gvy = gcy; } - //let gvy = gcy; } color_avg.0 /= count as f32; diff --git a/src/voxel/pipeline.rs b/src/voxel/pipeline.rs deleted file mode 100644 index fac322f..0000000 --- a/src/voxel/pipeline.rs +++ /dev/null @@ -1,799 +0,0 @@ -use std::num::NonZero; - -use bytemuck::cast_slice; -use crevice::std140::AsStd140; -use glam::Mat4; -use wgpu::BindGroup; -use wgpu::BindGroupDescriptor; -use wgpu::BindGroupEntry; -use wgpu::BindGroupLayout; -use wgpu::Buffer; -use wgpu::BufferUsages; -use wgpu::CommandEncoder; -use wgpu::CommandEncoderDescriptor; -use wgpu::ComputePass; -use wgpu::ComputePassDescriptor; -use wgpu::ComputePipeline; -use wgpu::Device; -use wgpu::Operations; -use wgpu::Queue; -use wgpu::RenderPipeline; -use wgpu::ShaderModuleDescriptor; -use wgpu::ShaderStages; -use wgpu::TextureFormat; -use wgpu::TextureView; -use wgpu::VertexBufferLayout; -use wgpu::util::StagingBelt; - -use crate::as_raw_bytes; -use crate::camera::Camera; -use crate::voxel::cache::ColorPoolElement; -use crate::voxel::cache::LocationPoolElement; -use crate::voxel::cache::RequestBufferElement; -use crate::voxel::cache::StructurePoolElement; -use crate::voxel::gpu::StructurePointer; -use crate::voxel::sparse::Color; - -// Represents a chunk to be rendered by the voxel pipeline -#[derive(Clone, Copy)] -#[repr(C)] -pub struct ChunkObject -{ - // Chunk object transform - pub transform: Mat4, - - // Chunk data - pub color: Color, - pub subdivided: bool, - - // Producer specific data - pub id: u32, -} - -#[derive(Clone, Copy)] -#[repr(C)] -pub struct CacheChunkObject -{ - // Chunk object transform - transform: Mat4, - - // Chunk data - color: Color, - - // Producer specific data - id: u32, - - // Pointer into cache - pointer: StructurePointer, -} - -#[derive(Clone, Copy)] -pub struct ChunkHandle(usize); - -pub struct CacheRequest -{ - // Records how many rays requested a - // resource - count: u32, -} - -pub struct VoxelPipeline -where - [(); N * N * N]:, -{ - chunk_allocations: Vec, - chunk_indices: Buffer, - chunk_staging: StagingBelt, - chunk_objects: Buffer, - chunk_requests: Buffer, - - chunk_objects_bind_group_layout: BindGroupLayout, - ray_bind_group_layout: BindGroupLayout, - chunk_objects_bind_group: BindGroup, - chunk_requests_bind_group: BindGroup, - render_pipeline: RenderPipeline, - - // Cache pools - structure_pool: Buffer, - color_pool: Buffer, - location_pool: Buffer, - - // Cache interaction - request_buffer: Buffer, - usage_buffer: Buffer, - cache_bind_group: BindGroup, - - device: Device, - queue: Queue, - - // Cache keeping shaders - clear_chunk_request: ComputePipeline, - sort_requests: ComputePipeline, -} - -#[derive(AsStd140)] -struct RenderPipelineImmediate -{ - view_proj: Mat4, -} - -impl VoxelPipeline -where - [(); N * N * N]:, -{ - pub fn new( - cache_size: usize, - device: Device, - queue: Queue, - surface_format: TextureFormat, - ) -> Self - { - let shader_module = device.create_shader_module(wgpu::ShaderModuleDescriptor { - label: Some("Main shader module"), - source: wgpu::ShaderSource::Wgsl( - std::fs::read_to_string("shaders/voxel.wgsl") - .unwrap() - .into(), - ), - }); - - let cache_bind_group_layout = - device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { - label: Some("Cache bind group layout"), - entries: &[ - // Location pool - wgpu::BindGroupLayoutEntry { - binding: 0, - visibility: ShaderStages::FRAGMENT, - ty: wgpu::BindingType::Buffer { - ty: wgpu::BufferBindingType::Storage { read_only: false }, - has_dynamic_offset: false, - min_binding_size: None, - }, - count: None, - }, - // Color pool - wgpu::BindGroupLayoutEntry { - binding: 1, - visibility: ShaderStages::FRAGMENT, - ty: wgpu::BindingType::Buffer { - ty: wgpu::BufferBindingType::Storage { read_only: false }, - has_dynamic_offset: false, - min_binding_size: None, - }, - count: None, - }, - // Location pool - wgpu::BindGroupLayoutEntry { - binding: 2, - visibility: ShaderStages::FRAGMENT, - ty: wgpu::BindingType::Buffer { - ty: wgpu::BufferBindingType::Storage { read_only: false }, - has_dynamic_offset: false, - min_binding_size: None, - }, - count: None, - }, - // Request buffer - wgpu::BindGroupLayoutEntry { - binding: 3, - visibility: ShaderStages::FRAGMENT, - ty: wgpu::BindingType::Buffer { - ty: wgpu::BufferBindingType::Storage { read_only: false }, - has_dynamic_offset: false, - min_binding_size: None, - }, - count: None, - }, - // Usage buffer - wgpu::BindGroupLayoutEntry { - binding: 4, - visibility: ShaderStages::FRAGMENT, - ty: wgpu::BindingType::Buffer { - ty: wgpu::BufferBindingType::Storage { read_only: false }, - has_dynamic_offset: false, - min_binding_size: None, - }, - count: None, - }, - ], - }); - - let chunk_objects_bind_group_layout = - device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { - label: Some("Chunk objects bg"), - entries: &[wgpu::BindGroupLayoutEntry { - binding: 0, - visibility: ShaderStages::VERTEX, - ty: wgpu::BindingType::Buffer { - ty: wgpu::BufferBindingType::Storage { read_only: true }, - has_dynamic_offset: false, - min_binding_size: None, - }, - count: None, - }], - }); - - let chunk_requests_bind_group_layout = - device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { - label: Some("Ray bind group layout"), - entries: &[wgpu::BindGroupLayoutEntry { - binding: 0, - visibility: ShaderStages::FRAGMENT | ShaderStages::COMPUTE, - ty: wgpu::BindingType::Buffer { - ty: wgpu::BufferBindingType::Storage { read_only: false }, - has_dynamic_offset: false, - min_binding_size: None, - }, - count: None, - }], - }); - - let request_buffer_sort_bind_group_layout = - device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { - label: Some("Ray bind group layout"), - entries: &[wgpu::BindGroupLayoutEntry { - binding: 0, - visibility: ShaderStages::COMPUTE, - ty: wgpu::BindingType::Buffer { - ty: wgpu::BufferBindingType::Storage { read_only: false }, - has_dynamic_offset: false, - min_binding_size: None, - }, - count: None, - }], - }); - - let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { - label: Some("Voxel pipeline layout"), - - bind_group_layouts: &[ - Some(&chunk_objects_bind_group_layout), - Some(&chunk_requests_bind_group_layout), - Some(&cache_bind_group_layout), - ], - immediate_size: RenderPipelineImmediate::std140_size_static() as u32, - }); - - let chunk_pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor { - label: Some("Render pipeline"), - layout: Some(&pipeline_layout), - vertex: wgpu::VertexState { - module: &shader_module, - entry_point: Some("chunk"), - compilation_options: Default::default(), - buffers: &[Some(VertexBufferLayout { - array_stride: size_of::() as u64, - step_mode: wgpu::VertexStepMode::Instance, - attributes: &[wgpu::VertexAttribute { - format: wgpu::VertexFormat::Uint32, - offset: 0, - shader_location: 0, - }], - })], - }, - primitive: wgpu::PrimitiveState { - topology: wgpu::PrimitiveTopology::TriangleList, - strip_index_format: None, - front_face: wgpu::FrontFace::Ccw, - cull_mode: None, - unclipped_depth: false, - polygon_mode: wgpu::PolygonMode::Fill, - conservative: false, - }, - depth_stencil: Some(wgpu::DepthStencilState { - format: wgpu::TextureFormat::Depth24PlusStencil8, - depth_write_enabled: Some(true), - depth_compare: Some(wgpu::CompareFunction::LessEqual), - stencil: wgpu::StencilState::default(), - bias: wgpu::DepthBiasState::default(), - }), - multisample: wgpu::MultisampleState::default(), - fragment: Some(wgpu::FragmentState { - module: &shader_module, - entry_point: Some("fragment"), - compilation_options: wgpu::PipelineCompilationOptions::default(), - targets: &[Some(wgpu::ColorTargetState { - format: surface_format, - blend: None, - write_mask: wgpu::ColorWrites::default(), - })], - }), - multiview_mask: None, - cache: None, - }); - - let chunk_indices = device.create_buffer(&wgpu::wgt::BufferDescriptor { - label: Some("Chunk index buffer"), - size: size_of::() as u64, - usage: BufferUsages::VERTEX | BufferUsages::COPY_DST, - mapped_at_creation: true, - }); - chunk_indices.unmap(); - - let chunk_objects = device.create_buffer(&wgpu::wgt::BufferDescriptor { - label: Some("Chunk buffer"), - size: size_of::() as u64, - usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC, - mapped_at_creation: false, - }); - - let chunk_requests = device.create_buffer(&wgpu::wgt::BufferDescriptor { - label: Some("Chunk request buffer"), - size: size_of::() as u64, - usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC, - mapped_at_creation: false, - }); - - // Pools - let structure_pool = device.create_buffer(&wgpu::wgt::BufferDescriptor { - label: Some("Structure pool"), - size: size_of::>() as u64 * cache_size as u64, - usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC, - mapped_at_creation: false, - }); - - let color_pool = device.create_buffer(&wgpu::wgt::BufferDescriptor { - label: Some("Color pool"), - size: size_of::>() as u64 * cache_size as u64, - usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC, - mapped_at_creation: false, - }); - - let location_pool = device.create_buffer(&wgpu::wgt::BufferDescriptor { - label: Some("Locatino pool"), - size: size_of::() as u64 * cache_size as u64, - usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC, - mapped_at_creation: false, - }); - - let request_buffer = device.create_buffer(&wgpu::wgt::BufferDescriptor { - label: Some("Request buffer"), - size: size_of::>() as u64 * cache_size as u64, - usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC, - mapped_at_creation: false, - }); - - let request_sort_buffer = device.create_buffer(&wgpu::wgt::BufferDescriptor { - label: Some("Request sort buffer"), - size: size_of::() as u64 * cache_size as u64, - usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC, - mapped_at_creation: true, - }); - - request_sort_buffer - .get_mapped_range_mut(0..) - .unwrap() - .copy_from_slice(cast_slice( - (0..cache_size as u32).collect::>().as_slice(), - )); - request_sort_buffer.unmap(); - - let usage_buffer = device.create_buffer(&wgpu::wgt::BufferDescriptor { - label: Some("Usage buffer buffer"), - size: size_of::() as u64 * cache_size as u64, - usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC, - mapped_at_creation: false, - }); - - let cache_bind_group = device.create_bind_group(&BindGroupDescriptor { - label: Some("Cache bind group"), - layout: &cache_bind_group_layout, - entries: &[ - // Structure pool - wgpu::BindGroupEntry { - binding: 0, - resource: structure_pool.as_entire_binding(), - }, - // Color pool - wgpu::BindGroupEntry { - binding: 1, - resource: color_pool.as_entire_binding(), - }, - // Location pool - wgpu::BindGroupEntry { - binding: 2, - resource: location_pool.as_entire_binding(), - }, - // Request buffer - wgpu::BindGroupEntry { - binding: 3, - resource: request_buffer.as_entire_binding(), - }, - // Usage buffer - wgpu::BindGroupEntry { - binding: 4, - resource: structure_pool.as_entire_binding(), - }, - ], - }); - - // Cache keeping shaders - let clear_chunk_request = - device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { - label: Some("Clean chunk request"), - layout: Some( - &device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { - label: Some("clear chunk requests layout"), - bind_group_layouts: &[Some(&chunk_requests_bind_group_layout)], - immediate_size: 0, - }), - ), - module: &device.create_shader_module(ShaderModuleDescriptor { - label: Some("clear_chunk_requests shader module"), - source: wgpu::ShaderSource::Wgsl( - " - @group(0) @binding(0) var chunk_requests: array; - @compute - @workgroup_size(16) - fn main( - @builtin(global_invocation_id) global_invocation_id: vec3 - ) - { - let index = global_invocation_id.x; - let total = arrayLength(&chunk_requests); - - if(index < total) - { - chunk_requests[index] = 0; - } - } - " - .into(), - ), - }), - entry_point: Some("main"), - compilation_options: Default::default(), - cache: None, - }); - - // Sorting requests - let sort_requests_bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor { - label: Some("Sort request bing group"), - layout: &request_buffer_sort_bind_group_layout, - entries: &[BindGroupEntry { - binding: 0, - resource: request_sort_buffer.as_entire_binding(), - }], - }); - let sort_requests = - device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { - label: Some("Sort node request"), - layout: Some( - &device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { - label: Some("Sort node requests"), - bind_group_layouts: &[Some(&cache_bind_group_layout), Some(&request_buffer_sort_bind_group_layout)], - immediate_size: 0, - }), - ), - module: &device.create_shader_module(ShaderModuleDescriptor { - label: Some("clear_chunk_requests shader module"), - source: wgpu::ShaderSource::Wgsl( - " - struct RequestElement - { - children: array, 64> - } - - @group(0) @binding(0) var structure_pool: array; - @group(0) @binding(1) var color_pool: array; - @group(0) @binding(2) var location_pool: array; - @group(0) @binding(3) var request_buffer: array; - @group(0) @binding(4) var usage_buffer: array; - - @group(1) @binding(0) var sort_indirection: array; - - @compute - @workgroup_size(16) - fn main( - @builtin(global_invocation_id) global_invocation_id: vec3 - ) - { - let index = global_invocation_id.x; - let total = arrayLength(&chunk_requests); - - // Odd pass - let a = index * 2 + 1; - let b = a + 1; - - // Gather elements - let av = request_buffer[sort_indirection[a]]; - let bv = request_buffer[sort_indirection[b]]; - - if b < total && av > bv - { - request_buffer[sort_indirection[a]] = bv; - request_buffer[sort_indirection[b]] = av; - } - storageBarrier(); - - // Even pass - a = index * 2; - b = a + 1; - - // Gather elements - av = request_buffer[sort_indirection[a]]; - bv = request_buffer[sort_indirection[b]]; - - if b < total && av > bv - { - request_buffer[sort_indirection[a]] = bv; - request_buffer[sort_indirection[b]] = av; - } - } - " - .into(), - ), - }), - entry_point: Some("main"), - compilation_options: Default::default(), - cache: None, - }); - - VoxelPipeline { - // Only one slot, no allocated chunks at the beginning - chunk_allocations: vec![false], - chunk_objects_bind_group: device.create_bind_group(&wgpu::BindGroupDescriptor { - label: Some("Chunk objects bind group"), - layout: &chunk_objects_bind_group_layout, - entries: &[wgpu::BindGroupEntry { - binding: 0, - resource: chunk_objects.as_entire_binding(), - }], - }), - chunk_requests_bind_group: device.create_bind_group(&wgpu::BindGroupDescriptor { - label: Some("Ray bind group"), - layout: &chunk_requests_bind_group_layout, - entries: &[wgpu::BindGroupEntry { - binding: 0, - resource: chunk_requests.as_entire_binding(), - }], - }), - chunk_objects, - chunk_requests, - chunk_indices, - - structure_pool, - color_pool, - location_pool, - request_buffer, - usage_buffer, - cache_bind_group, - - chunk_objects_bind_group_layout, - ray_bind_group_layout: chunk_requests_bind_group_layout, - chunk_staging: StagingBelt::new(device.clone(), size_of::() as u64), - render_pipeline: chunk_pipeline, - device, - queue, - - clear_chunk_request, - sort_requests, - } - } - - pub fn render( - &mut self, - encoder: &mut CommandEncoder, - texture_view: &TextureView, - depth_buffer_view: &TextureView, - camera: &Camera, - ) - { - let mut renderpass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { - label: None, - color_attachments: &[Some(wgpu::RenderPassColorAttachment { - view: texture_view, - depth_slice: None, - resolve_target: None, - ops: wgpu::Operations { - load: wgpu::LoadOp::Clear(wgpu::Color::BLACK), - store: wgpu::StoreOp::Store, - }, - })], - depth_stencil_attachment: Some(wgpu::RenderPassDepthStencilAttachment { - view: depth_buffer_view, - depth_ops: Some(Operations { - load: wgpu::LoadOp::Clear(1.), - store: wgpu::StoreOp::Discard, - }), - stencil_ops: None, - }), - timestamp_writes: None, - occlusion_query_set: None, - multiview_mask: None, - }); - - renderpass.set_pipeline(&self.render_pipeline); - renderpass.set_bind_group(0, Some(&self.chunk_objects_bind_group), &[]); - renderpass.set_bind_group(1, Some(&self.chunk_requests_bind_group), &[]); - renderpass.set_bind_group(2, Some(&self.cache_bind_group), &[]); - renderpass.set_vertex_buffer(0, self.chunk_indices.slice(0..)); - renderpass.set_immediates( - 0, - RenderPipelineImmediate { - view_proj: camera.view_proj(), - } - .as_std140() - .as_bytes(), - ); - renderpass.draw( - 0..36, - 0..(self.chunk_allocations.iter().filter(|x| **x).count() as u32), - ); - - // End the renderpass. - drop(renderpass); - - let mut compute_pass = encoder.begin_compute_pass(&ComputePassDescriptor { - label: Some("cache keeping pass"), - timestamp_writes: None, - }); - - compute_pass.set_bind_group(0, Some(&self.chunk_requests_bind_group), &[]); - compute_pass.set_pipeline(&self.clear_chunk_request); - compute_pass.dispatch_workgroups( - self.chunk_allocations.len().next_multiple_of(16) as u32 / 16, - 1, - 1, - ); - - drop(compute_pass) - } - - fn update_indices(&mut self) - { - let indices = self - .chunk_allocations - .iter() - .enumerate() - .filter(|(_, b)| **b) - .map(|(i, _)| i as u32) - .collect::>(); - - self.chunk_indices = self.device.create_buffer(&wgpu::wgt::BufferDescriptor { - label: Some("Chunk index buffer"), - size: size_of::() as u64 * indices.len() as u64, - usage: BufferUsages::VERTEX | BufferUsages::COPY_DST, - mapped_at_creation: true, - }); - - // Copy chunk indices into new buffer - self.chunk_indices - .get_mapped_range_mut(0..(size_of::() as u64 * indices.len() as u64)) - .unwrap() - .copy_from_slice(cast_slice(&indices)); - self.chunk_indices.unmap(); - } - - pub fn remove_chunks(&mut self, handles: &[ChunkHandle]) - { - for handle in handles.iter() - { - self.chunk_allocations[handle.0] = false; - } - self.update_indices(); - } - - pub fn push_new_chunks(&mut self, objects: &[ChunkObject]) -> Vec - { - // Find room for new chunks - let mut destinations = vec![0; objects.len()]; - - let mut encoder = self - .device - .create_command_encoder(&CommandEncoderDescriptor { - label: Some("Chunk buffer writes"), - }); - - // count available space - let space = self.chunk_allocations.iter().filter(|x| !*x).count(); - if space < objects.len() - { - // Allocate more space - // Get first bigger power of two - let necessary_space = - (objects.len() + self.chunk_allocations.len()).next_power_of_two(); - dbg!(necessary_space); - self.chunk_allocations - .extend(vec![false; necessary_space - self.chunk_allocations.len()]); - - // Make buffer bigger - let old_buffer = self.chunk_objects.clone(); - self.chunk_objects = self.device.create_buffer(&wgpu::wgt::BufferDescriptor { - label: Some("Chunk buffer"), - size: size_of::() as u64 * self.chunk_allocations.len() as u64, - usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC, - mapped_at_creation: false, - }); - - let old_chunk_requests = self.chunk_requests.clone(); - self.chunk_requests = self.device.create_buffer(&wgpu::wgt::BufferDescriptor { - label: Some("Chunk request buffer"), - size: size_of::() as u64 * self.chunk_allocations.len() as u64, - usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC, - mapped_at_creation: false, - }); - - self.chunk_objects_bind_group = - self.device.create_bind_group(&wgpu::BindGroupDescriptor { - label: Some("Chunk objects bind group"), - layout: &self.chunk_objects_bind_group_layout, - entries: &[wgpu::BindGroupEntry { - binding: 0, - resource: self.chunk_objects.as_entire_binding(), - }], - }); - - self.chunk_requests_bind_group = - self.device.create_bind_group(&wgpu::BindGroupDescriptor { - label: Some("Ray bind group"), - layout: &self.ray_bind_group_layout, - entries: &[wgpu::BindGroupEntry { - binding: 0, - resource: self.chunk_requests.as_entire_binding(), - }], - }); - - encoder.copy_buffer_to_buffer( - &old_chunk_requests, - 0, - &self.chunk_requests, - 0, - old_chunk_requests.size(), - ); - - encoder.copy_buffer_to_buffer( - &old_buffer, - 0, - &self.chunk_objects, - 0, - old_buffer.size(), - ); - } - - let mut dest_ptr = 0; - for (i, allocated) in self - .chunk_allocations - .iter_mut() - .enumerate() - .filter(|(_, allocated)| !**allocated) - .take(destinations.len()) - { - if !*allocated - { - *allocated = true; - destinations[dest_ptr] = i; - dest_ptr += 1; - } - } - - // Write each new chunk - for (destination, object) in destinations.iter().zip(objects.iter()) - { - let cache_object = CacheChunkObject { - transform: object.transform, - id: object.id, - color: object.color, - pointer: StructurePointer::new(object.subdivided, false, 0), - }; - - let mut view = self.chunk_staging.write_buffer( - &mut encoder, - &self.chunk_objects, - *destination as u64 * size_of::() as u64, - NonZero::new(size_of::() as u64).unwrap(), - ); - - let temp_slice = [cache_object]; - view.copy_from_slice(unsafe { as_raw_bytes(&temp_slice) }); - } - - self.update_indices(); - - self.chunk_staging.finish_and_recall_on_submit(&encoder); - self.queue.submit([encoder.finish()]); - - destinations.iter().map(|d| ChunkHandle(*d)).collect() - } -} diff --git a/src/voxel.rs b/src/voxel_cache.rs similarity index 100% rename from src/voxel.rs rename to src/voxel_cache.rs diff --git a/src/voxel/cache.rs b/src/voxel_cache/cache.rs similarity index 89% rename from src/voxel/cache.rs rename to src/voxel_cache/cache.rs index cf5ca35..a2a0742 100644 --- a/src/voxel/cache.rs +++ b/src/voxel_cache/cache.rs @@ -13,10 +13,11 @@ use wgpu::ComputePipeline; use wgpu::Device; use wgpu::Queue; use wgpu::ShaderStages; +use wgpu::util::DeviceExt; use wgpu::util::StagingBelt; -use crate::voxel::gpu::StructurePointer; -use crate::voxel::sparse::Color; +use crate::voxel_cache::gpu::StructurePointer; +use crate::voxel_cache::sparse::Color; #[derive(Clone, Copy, Zeroable, Pod, Debug)] #[repr(C)] @@ -36,6 +37,9 @@ pub struct RequestBuffer pub cache_size: usize, pub request_buffer: Buffer, pub request_count_buffer: Buffer, + pub indirect_count_storage: Buffer, + pub element_sort_count_buffer: Buffer, + pub indirect_count: Buffer, pub sort_buffer: Buffer, pub reset_pipeline: ComputePipeline, pub sort_pipeline: ComputePipeline, @@ -74,6 +78,28 @@ where mapped_at_creation: false, }); + + let indirect_count = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("indirect_count_buffer"), + usage: BufferUsages::INDIRECT | BufferUsages::COPY_DST, + size: size_of::() as u64, + mapped_at_creation: false, + }); + + let indirect_count_storage = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("indirect_count_storage_buffer"), + usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC, + size: size_of::() as u64, + mapped_at_creation: false, + }); + + let element_sort_count_buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("indirect_count_storage_buffer"), + usage: BufferUsages::STORAGE, + size: size_of::() as u64, + mapped_at_creation: false, + }); + // One element number per child entry // One number to identify node, one to identify sub child let sort_buffer_count = cache_size * N * N * N; @@ -128,6 +154,26 @@ where }, count: None, }, + wgpu::BindGroupLayoutEntry { + binding: 3, + visibility: ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: false }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 4, + visibility: ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: false }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, ], }); @@ -147,6 +193,14 @@ where binding: 2, resource: request_count_buffer.as_entire_binding(), }, + BindGroupEntry { + binding: 3, + resource: indirect_count_storage.as_entire_binding(), + }, + BindGroupEntry { + binding: 4, + resource: element_sort_count_buffer.as_entire_binding(), + }, ], }); @@ -181,6 +235,8 @@ where @group(0) @binding(0) var request_buffer: array; @group(0) @binding(1) var sort_indirection: array; @group(0) @binding(2) var request_count: u32; + @group(0) @binding(3) var indirect_count: vec3; + @group(0) @binding(4) var element_sort_count: u32; fn big_fusion(index: u32, phase_size: u32) -> vec2 {{ @@ -218,30 +274,24 @@ where let index = global_invocation_id.x; let total = arrayLength(&sort_indirection); - /* - let request_count_round_up = - ((request_count / {children_count}) + select(u32(0), u32(1), request_count % {children_count} != 0)) * {children_count}; - if(index * 2 >= request_count_round_up) - {{ return; }} - */ - let sub_phase = (both_phase >> 16) & 0xFFFF; let phase = both_phase & 0xFFFF; let phase_total_width = u32(1 << (phase + 1)); - request_count = 0; - // Check if phase is last /* - let phase_count = u32(ceil(log2(f32(request_count)))); - if(phase > phase_count) + if(phase == sub_phase) {{ - return; + let next_phase_width = phase_total_width * 2; + if(next_phase_width >= element_sort_count && phase_total_width >= element_sort_count) + {{ + // This was the final phase, stop + indirect_count = vec3(0); + }} }} */ - /* if(phase == 0) {{ let a = index * 2; @@ -249,7 +299,6 @@ where sort_indirection[a].child = a % {children_count}; sort_indirection[b].child = b % {children_count}; }} - */ var swap_indices = vec2(0, 0); @@ -265,7 +314,7 @@ where // Do swap //if(swap_indices.y >= request_count_round_up) - if(swap_indices.y >= total) + if(swap_indices.y >= element_sort_count) {{ // Suppose that swap_indices.y is -inf, dont swap return; @@ -313,7 +362,9 @@ where @group(0) @binding(0) var request_buffer: array; @group(0) @binding(1) var sort_indirection: array; - @group(0) @binding(2) var request_counts: u32; + @group(0) @binding(2) var request_counts: atomic; + @group(0) @binding(3) var indirect_count: u32; + @group(0) @binding(4) var node_sort_count: u32; @compute @workgroup_size(64) @@ -321,57 +372,55 @@ where @builtin(global_invocation_id) global_invocation_id: vec3 ) {{ - let index = global_invocation_id.x; - let sindex = index * {children_count}; - let len = arrayLength(&request_buffer); - if(index >= len) - {{ + let index = global_invocation_id.x; + let len = arrayLength(&request_buffer); + if(index > len) {{ return; }} + let sindex = index * {children_count}; + + if(phase == 0) + {{ + var count = 0; + for(var i = 0; i < {children_count}; i++) + {{ + count += select(0, 1, request_buffer[index].children[i] != 0); + }} + + sort_indirection[sindex].child = select(u32(0), u32(1), count != 0); + if count != 0 + {{ + atomicAdd(&request_counts, 1); + }} return; - }} + }} - if(phase == u32(0xFFFFFFFF)) - {{ - request_counts = 0; - var children_sum = u32(0); - for(var i = u32(0); i < {children_count}; i++) - {{ - children_sum += request_buffer[index].children[i]; - }} - sort_indirection[sindex].child = select(u32(0), u32(1), children_sum != 0); - return; - }} + if(phase == 0xFFFFFFFF) + {{ + // Double buffering bring back + sort_indirection[sindex].child = sort_indirection[sindex].node; + }} - if(phase == u32(0x0FFFFFFF)) - {{ - sort_indirection[sindex].child = sort_indirection[sindex].node; - return; - }} - - let phase_offset = u32(1) << phase; - let gather_offset = phase_offset * u32({children_count}); - if(index >= phase_offset) - {{ - // Double buffering - if(phase % 2 == 0) - {{ - sort_indirection[sindex].node = sort_indirection[sindex].child + sort_indirection[sindex - gather_offset].child; - }}else - {{ - sort_indirection[sindex].child = sort_indirection[sindex].node + sort_indirection[sindex - gather_offset].node; - }} - }} - else - {{ - // Double buffering - if(phase % 2 == 0) - {{ - sort_indirection[sindex].node = sort_indirection[sindex].child; - }}else - {{ - sort_indirection[sindex].child = sort_indirection[sindex].node; - }} - }} + // Phase is not zero, running sum part + let running_sum_phase = phase - 1; + let running_sum_offset = u32((1 << running_sum_phase) * {children_count}); + var add = u32(0); + // Double buffering + if(running_sum_phase % 2 == 0) + {{ + if(sindex >= running_sum_offset) + {{ + add = sort_indirection[sindex - running_sum_offset].child; + }} + sort_indirection[sindex].node = sort_indirection[sindex].child + add; + }} + else + {{ + if(sindex >= running_sum_offset) + {{ + add = sort_indirection[sindex - running_sum_offset].node; + }} + sort_indirection[sindex].child = sort_indirection[sindex].node + add; + }} }} ") .into(), @@ -406,6 +455,8 @@ where @group(0) @binding(0) var request_buffer: array; @group(0) @binding(1) var sort_indirection: array; @group(0) @binding(2) var request_counts: atomic; + @group(0) @binding(3) var indirect_count: vec3; + @group(0) @binding(4) var element_sort_count: u32; @compute @workgroup_size(64) @@ -416,27 +467,28 @@ where let index = global_invocation_id.x; let sindex = index * {children_count}; let len = arrayLength(&request_buffer); - if(index >= len) + + element_sort_count = sort_indirection[(len - 1) * {children_count}].child * {children_count}; + let invocation_count = (element_sort_count / 2) + select(u32(0), u32(1), element_sort_count % 2 != 0); + indirect_count = vec3( + (invocation_count / 64) + select(u32(0), u32(1), invocation_count % 64 != 0), + 1, 1 + ); + let destination_index = sort_indirection[sindex].child - 1; + + var count = u32(0); + for(var i = 0; i < {children_count}; i++) {{ - return; + count += request_buffer[index].children[i]; }} - // Check if keep - var node_request_count = u32(0); - for(var i = 0; i < {children_count}; i++) - {{ - node_request_count += select(u32(0), u32(1), request_buffer[index].children[i] != 0); - }} - let keep = node_request_count != 0; - atomicAdd(&request_counts, node_request_count); - if(keep) - {{ - let running_sum_value = sort_indirection[sindex].child; - for(var i = u32(0); i < {children_count}; i++) + if(count != 0) // Keep ? + {{ + for(var i = u32(0); i < u32({children_count}); i++) {{ - sort_indirection[running_sum_value - 1 + i].node = index; + sort_indirection[destination_index * u32({children_count}) + i].node = index; }} - }} + }} }} ") .into(), @@ -462,7 +514,8 @@ where @group(0) @binding(0) var request_buffer: array; @group(0) @binding(1) var _ignore: array; - @group(0) @binding(2) var request_count: atomic; + @group(0) @binding(2) var request_count: u32; + @group(0) @binding(4) var node_sort_count: u32; @compute @workgroup_size(64) @@ -479,12 +532,15 @@ where {{ for(var i = 0; i < {children_count}; i += 1) {{ + /* if(request_buffer[index].children[i] != 0) {{ atomicAdd(&request_count, 1); }} + */ request_buffer[index].children[i] = 0; }} + request_count = 0; }} }} ") @@ -500,6 +556,9 @@ where cache_size, request_buffer, request_count_buffer, + indirect_count, + indirect_count_storage, + element_sort_count_buffer, running_sum_pipeline, sort_buffer, sort_pipeline, @@ -542,58 +601,71 @@ where pub fn sort_requests(&self, encoder: &mut CommandEncoder) { - let mut compute_pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { - label: Some("request_buffer_sort_compute_pass"), + let mut compaction_compute_pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: Some("request_buffer_compaction_compute_pass"), timestamp_writes: None, }); - compute_pass.set_bind_group(0, Some(&self.bindgroup), &[]); + compaction_compute_pass.set_bind_group(0, Some(&self.bindgroup), &[]); let request_element_count = self.cache_size; // Each child slot is sorted let workgroups_invocations = request_element_count.div_ceil(64); - /* // = Perform list compaction // == Running sum - compute_pass.set_pipeline(&self.running_sum_pipeline); + compaction_compute_pass.set_pipeline(&self.running_sum_pipeline); // Phase 0: Put ones in correct location - compute_pass.set_immediates(0, bytes_of(&u32::MAX)); - compute_pass.dispatch_workgroups(workgroups_invocations as u32, 1, 1); + compaction_compute_pass.set_immediates(0, bytes_of(&0)); + compaction_compute_pass.dispatch_workgroups(workgroups_invocations as u32, 1, 1); // Phase _: running sum // Running sum phase - for i in 0..(request_element_count.next_power_of_two().ilog2()) + let running_sum_steps = request_element_count.next_power_of_two().ilog2(); + for i in 1..=running_sum_steps { - compute_pass.set_immediates(0, bytes_of(&i)); - compute_pass.dispatch_workgroups(workgroups_invocations as u32, 1, 1); + compaction_compute_pass.set_immediates(0, bytes_of(&i)); + compaction_compute_pass.dispatch_workgroups(workgroups_invocations as u32, 1, 1); } - if !request_element_count.next_power_of_two().ilog2().is_multiple_of(2) + if !running_sum_steps.is_multiple_of(2) { - compute_pass.set_immediates(0, bytes_of(&0x0FFFFFFF)); - compute_pass.dispatch_workgroups(workgroups_invocations as u32, 1, 1); + // Bring back double buffer + compaction_compute_pass.set_immediates(0, bytes_of(&0xFFFFFFFF_u32)); + compaction_compute_pass.dispatch_workgroups(workgroups_invocations as u32, 1, 1); } // == Stream compaction - compute_pass.set_pipeline(&self.compaction_pipeline); - compute_pass.dispatch_workgroups(workgroups_invocations as u32, 1, 1); - */ + compaction_compute_pass.set_pipeline(&self.compaction_pipeline); + compaction_compute_pass.dispatch_workgroups(workgroups_invocations as u32, 1, 1); + drop(compaction_compute_pass); + + // == Copy count into indirect buffer + + + // = Sort // == Bitonic sort - compute_pass.set_pipeline(&self.sort_pipeline); let sort_element_count = self.cache_size * N * N * N; let phases_upper_bound = sort_element_count.next_power_of_two().ilog2(); // Phase 0 dispatch all for i in 0..=phases_upper_bound { + encoder.copy_buffer_to_buffer(&self.indirect_count_storage, 0, &self.indirect_count, 0, Some(size_of::() as u64)); + let mut sort_compute_pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: Some(format!("request_buffer_sort_compute_pass_{}", 0).as_str()), + timestamp_writes: None, + }); + sort_compute_pass.set_bind_group(0, Some(&self.bindgroup), &[]); + sort_compute_pass.set_pipeline(&self.sort_pipeline); for j in 0..=i { - compute_pass.set_immediates(0, bytemuck::bytes_of( + sort_compute_pass.set_immediates(0, bytemuck::bytes_of( &(i | (j << 16)) )); - compute_pass.dispatch_workgroups(sort_element_count.div_ceil(2).div_ceil(64) as u32, 1, 1); + sort_compute_pass.dispatch_workgroups_indirect(&self.indirect_count, 0); } + drop(sort_compute_pass); } } } @@ -1631,7 +1703,7 @@ where }} let overwritten_element = lru_list[index]; - if(parameters.write_pointers == 0) + if(parameters.write_pointers == 0 && usage_buffer[overwritten_element] != parameters.frame_timestamp) {{ // Phase 1 // Copy into cache page @@ -1645,7 +1717,8 @@ where // Mark dirty/correct timestamp usage_buffer[overwritten_element] = parameters.frame_timestamp + 1; }} - else + + if(parameters.write_pointers != 0 && usage_buffer[overwritten_element] != parameters.frame_timestamp) {{ // Phase 2 @@ -1863,14 +1936,10 @@ where // Sort requests, reset request buffers to count requests self.usage_buffer.sort_usage(encoder); self.request_buffer.sort_requests(encoder); - self.request_buffer.reset_requests(encoder); self.structure_table .request_buffer_mut() .sort_requests(encoder); - self.structure_table - .request_buffer_mut() - .reset_requests(encoder); // Sorted requests are now in the group let user_target_bindgroup = self.device.create_bind_group(&wgpu::BindGroupDescriptor { @@ -1899,6 +1968,12 @@ where let shader_invocation_count = request_target.size() / size_of::() as u64; let workgroup_invocation_count = shader_invocation_count.div_ceil(64); write_requests_pass.dispatch_workgroups(workgroup_invocation_count as u32, 1, 1); + drop(write_requests_pass); + + self.request_buffer.reset_requests(encoder); + self.structure_table + .request_buffer_mut() + .reset_requests(encoder); } pub fn current_timestamp(&self) -> u32 diff --git a/src/voxel/cache.wait.rs b/src/voxel_cache/cache.wait.rs similarity index 100% rename from src/voxel/cache.wait.rs rename to src/voxel_cache/cache.wait.rs diff --git a/src/voxel/gpu.rs b/src/voxel_cache/gpu.rs similarity index 93% rename from src/voxel/gpu.rs rename to src/voxel_cache/gpu.rs index a9e6811..827267a 100644 --- a/src/voxel/gpu.rs +++ b/src/voxel_cache/gpu.rs @@ -1,7 +1,7 @@ use bytemuck::Pod; use bytemuck::Zeroable; -use crate::voxel::sparse::Color; +use crate::voxel_cache::sparse::Color; #[derive(Clone, Copy, Pod, Zeroable)] #[repr(transparent)] diff --git a/src/voxel_cache/pipeline.rs b/src/voxel_cache/pipeline.rs new file mode 100644 index 0000000..2ca03f2 --- /dev/null +++ b/src/voxel_cache/pipeline.rs @@ -0,0 +1,68 @@ +use std::num::NonZero; + +use bytemuck::cast_slice; +use crevice::std140::AsStd140; +use glam::Mat4; +use wgpu::BindGroup; +use wgpu::BindGroupDescriptor; +use wgpu::BindGroupEntry; +use wgpu::BindGroupLayout; +use wgpu::Buffer; +use wgpu::BufferUsages; +use wgpu::CommandEncoder; +use wgpu::CommandEncoderDescriptor; +use wgpu::ComputePass; +use wgpu::ComputePassDescriptor; +use wgpu::ComputePipeline; +use wgpu::Device; +use wgpu::Operations; +use wgpu::Queue; +use wgpu::RenderPipeline; +use wgpu::ShaderModuleDescriptor; +use wgpu::ShaderStages; +use wgpu::TextureFormat; +use wgpu::TextureView; +use wgpu::VertexBufferLayout; +use wgpu::util::StagingBelt; + +use crate::as_raw_bytes; +use crate::camera::Camera; +use crate::voxel_cache::cache::ColorPoolElement; +use crate::voxel_cache::cache::LocationPoolElement; +use crate::voxel_cache::cache::RequestBufferElement; +use crate::voxel_cache::cache::StructurePoolElement; +use crate::voxel_cache::gpu::StructurePointer; +use crate::voxel_cache::sparse::Color; + +// Represents a chunk to be rendered by the voxel pipeline +#[derive(Clone, Copy)] +#[repr(C)] +pub struct ChunkObject +{ + // Chunk object transform + pub transform: Mat4, + + // Chunk data + pub color: Color, + pub subdivided: bool, + + // Producer specific data + pub id: u32, +} + +#[derive(Clone, Copy)] +#[repr(C)] +pub struct CacheChunkObject +{ + // Chunk object transform + transform: Mat4, + + // Chunk data + color: Color, + + // Producer specific data + id: u32, + + // Pointer into cache + pointer: StructurePointer, +} diff --git a/src/voxel_cache/request_buffer.rs b/src/voxel_cache/request_buffer.rs new file mode 100644 index 0000000..e69de29 diff --git a/src/voxel/sparse.rs b/src/voxel_cache/sparse.rs similarity index 99% rename from src/voxel/sparse.rs rename to src/voxel_cache/sparse.rs index c76a69c..69b1aea 100644 --- a/src/voxel/sparse.rs +++ b/src/voxel_cache/sparse.rs @@ -5,8 +5,8 @@ use bytemuck::Pod; use bytemuck::Zeroable; use itertools::Itertools; -use crate::voxel::gpu::ExplicitNTreeNode; -use crate::voxel::gpu::StructurePointer; +use crate::voxel_cache::gpu::ExplicitNTreeNode; +use crate::voxel_cache::gpu::StructurePointer; #[derive(Debug, Clone, Copy, PartialEq, Pod, Zeroable)] #[repr(C)] diff --git a/src/voxel_cache/structure_table.rs b/src/voxel_cache/structure_table.rs new file mode 100644 index 0000000..e69de29 diff --git a/src/voxel_cache/usage_buffer.rs b/src/voxel_cache/usage_buffer.rs new file mode 100644 index 0000000..e69de29