From 2fbce4ab81c4a3ab3af7750cb58d8617fd776545 Mon Sep 17 00:00:00 2001 From: Albin Chaboissier Date: Wed, 2 Sep 2026 15:39:25 +0200 Subject: [PATCH] Adding histogram --- Cargo.toml | 1 + src/main.rs | 54 ++++++++++++++++++++--- src/voxel/cache.rs | 107 +++++++++++++++++++++++++++++---------------- 3 files changed, 118 insertions(+), 44 deletions(-) diff --git a/Cargo.toml b/Cargo.toml index 49b5d49..d0c98c6 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -9,6 +9,7 @@ crevice = {version = "0.20.1", features = ["glam"]} egui = "0.36.1" egui-wgpu = "0.36.1" egui-winit = "0.36.1" +egui_plot = "0.37.0" env_logger = "0.11.11" glam = "0.33.5" itertools = "0.15.0" diff --git a/src/main.rs b/src/main.rs index 54782d0..ac2eb9b 100644 --- a/src/main.rs +++ b/src/main.rs @@ -15,6 +15,7 @@ use crevice::std140::AsStd140; use crevice::std430::AsStd430; use egui::emath::fast_midpoint; use egui::mutex::Mutex; +use egui_plot::BarChart; use glam::Mat4; use glam::Vec3; use glam::Vec4; @@ -102,6 +103,7 @@ struct State pipeline: RenderPipeline, voxel_cache: Arc>>, + usage_vec: Arc>>, insertion_debounce: bool, camera: Camera, @@ -158,7 +160,7 @@ impl State let egui_renderer = EguiRenderer::new(&device, surface_format, &window); - let mut voxel_cache = VoxelCache::<4>::new(4096, device.clone(), queue.clone()); + let mut voxel_cache = VoxelCache::<4>::new(16000, device.clone(), queue.clone()); let id = voxel_cache.structure_table.allocate_structure(true); println!("id: {id}"); @@ -228,6 +230,7 @@ impl State depth_buffer: Self::create_depth_buffer(&device, size.width, size.height), queue, device, + usage_vec: Arc::new(Mutex::new(vec![])), pipeline: chunk_pipeline, voxel_cache: Arc::new(Mutex::new(voxel_cache)), insertion_debounce: false, @@ -407,7 +410,7 @@ impl State let requests = self.device.create_buffer(&wgpu::BufferDescriptor { label: Some("dummy_dumb_dinky_aaaahhh_buffer"), - size: 16 * 64, + size: 16 * 256, usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC, mapped_at_creation: false, }); @@ -419,9 +422,23 @@ impl State // If you wanted to call any drawing commands, they would go here. { self.egui_renderer.begin_frame(&self.window); - egui::Window::new("Window ! ") - .resizable(true) - .show(self.egui_renderer.context(), |ui| ui.label("Hello !")); + egui::Window::new("Window ! ").resizable(true).show( + self.egui_renderer.context(), + |ui| { + egui_plot::Plot::new("Plot").show(ui, |plot_ui| { + plot_ui.bar_chart(BarChart::new( + "histo", + self.usage_vec + .lock() + .iter() + .enumerate() + .map(|(i, x)| egui_plot::Bar::new(i as f64, *x as f64)) + .collect(), + )); + }); + ui.label("Hello !") + }, + ); let screen_descriptor = egui_wgpu::ScreenDescriptor { size_in_pixels: [self.size.width, self.size.height], @@ -438,6 +455,26 @@ impl State ); } + // Report usage + let time_stamp = self.voxel_cache.lock().current_timestamp(); + let cloned_usage_histogram = self.usage_vec.clone(); + DownloadBuffer::read_buffer( + &self.device, + &self.queue, + &self.voxel_cache.lock().usage_buffer.usage_buffer.slice(..), + move |slice| { + let slice = slice.unwrap(); + let cache_slice: &[u32] = bytemuck::cast_slice(&slice); + let mut v = vec![0; 100]; + for x in cache_slice.iter() + { + let diff = (time_stamp - x).min(99); + v[diff as usize] += 1; + } + *cloned_usage_histogram.lock() = v; + }, + ); + // Submit the command in the queue to execute self.queue.submit([encoder.finish()]); self.window.pre_present_notify(); @@ -446,14 +483,16 @@ impl State .camera .pressed_keyset .contains(&winit::keyboard::KeyCode::KeyF) + && self.insertion_debounce { self.insertion_debounce = false; } - if (!self + if (self .camera .pressed_keyset .contains(&winit::keyboard::KeyCode::KeyF)) + && !self.insertion_debounce { self.insertion_debounce = true; let request_count = self.voxel_cache.lock().total_request_count(); @@ -566,7 +605,6 @@ impl State ); cloned_queue.submit([encoder.finish()]); } - cloned_cache.lock().next_frame(); }, ); println!("Total request count: {}", request_count); @@ -578,6 +616,8 @@ impl State } let _ = self.device.poll(wgpu::wgt::PollType::Poll); } + + self.voxel_cache.lock().next_frame(); drop(rx); } } diff --git a/src/voxel/cache.rs b/src/voxel/cache.rs index 226cc53..d96ea10 100644 --- a/src/voxel/cache.rs +++ b/src/voxel/cache.rs @@ -601,7 +601,7 @@ impl UsageBuffer let usage_buffer = device.create_buffer(&wgpu::BufferDescriptor { label: Some("Usage buffer"), size: (size_of::() * cache_size) as u64, - usage: BufferUsages::STORAGE | BufferUsages::COPY_DST, + usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC, mapped_at_creation: true, }); let sort_buffer = device.create_buffer(&wgpu::BufferDescriptor { @@ -671,7 +671,7 @@ impl UsageBuffer let pipeline_layouts = &device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { label: Some("usage_buffer_pipeline_layout"), bind_group_layouts: &[Some(&bindgroup_layout)], - immediate_size: 0, + immediate_size: size_of::() as u32, }); let sort_pipeline = @@ -682,8 +682,36 @@ impl UsageBuffer label: Some("clear_chunk_requests shader module"), source: wgpu::ShaderSource::Wgsl( " - @group(0) @binding(0) var usage_stamps: array; - @group(0) @binding(1) var sort_indirection: array; + + @group(0) @binding(0) var usage_buffer: array; + @group(0) @binding(1) var sort_buffer: array; + var both_phase: u32; + + fn big_fusion(index: u32, phase_size: u32) -> vec2 + {{ + // Find out in which block index this invocation pertains + let block_index = index / (phase_size / 2); + let element_index = index % (phase_size / 2); + + let offset = block_index * phase_size; + let a = offset + element_index; + let b = offset + phase_size - 1 - element_index; + return vec2(a, b); + }} + + fn small_fusion(index: u32, phase: u32, sub_phase: u32) -> vec2 + {{ + let phase_size = u32(1 << (phase - sub_phase + 1)); + + let block_index = index / (phase_size / 2); + let element_index = index % (phase_size / 2); + + let offset = block_index * phase_size; + let a = offset + element_index; + let b = a + (phase_size / 2); + return vec2(a, b); + }} + @compute @workgroup_size(16) @@ -691,39 +719,40 @@ impl UsageBuffer @builtin(global_invocation_id) global_invocation_id: vec3 ) {{ - let index = global_invocation_id.x; - let total = arrayLength(&sort_indirection); - // Odd pass - var a = index * 2 + 1; - var b = a + 1; + let length = arrayLength(&sort_buffer); + let index = global_invocation_id.x; + let sub_phase = (both_phase >> 16) & 0xFFFF; + let phase = both_phase & 0xFFFF; + let phase_total_width = u32(1 << (phase + 1)); - // Gather elements - var av = usage_stamps[sort_indirection[a]]; - var bv = usage_stamps[sort_indirection[b]]; + var swap_indices = vec2(0, 0); - if b < total && av > bv // Reverse sort order - {{ - let temp = sort_indirection[a]; - sort_indirection[a] = sort_indirection[b]; - sort_indirection[b] = temp; - }} - storageBarrier(); + if(sub_phase == 0) + {{ + // Bitonic fusion + swap_indices = big_fusion(index, phase_total_width); + }}else + {{ + swap_indices = small_fusion(index, phase, sub_phase); + }} - // Even pass - a = index * 2; - b = a + 1; + // Do swap + if(swap_indices.y >= length) + {{ + // Suppose that swap_indices.y is -inf, dont swap + return; + }} + + let av = usage_buffer[sort_buffer[swap_indices.x]]; + let bv = usage_buffer[sort_buffer[swap_indices.y]]; - // Gather elements - av = usage_stamps[sort_indirection[a]]; - bv = usage_stamps[sort_indirection[b]]; - - if b < total && av > bv // Reverse sort order - {{ - let temp = sort_indirection[a]; - sort_indirection[a] = sort_indirection[b]; - sort_indirection[b] = temp; - }} + if(bv < av) + {{ + let temp = sort_buffer[swap_indices.x]; + sort_buffer[swap_indices.x] = sort_buffer[swap_indices.y]; + sort_buffer[swap_indices.y] = temp; + }} }} " .into(), @@ -777,15 +806,19 @@ impl UsageBuffer compute_pass.set_pipeline(&self.sort_pipeline); // Compute required shader invocations - let sort_element_count = self.cache_size; // Each element is sorted - let shader_invocations = sort_element_count.div_ceil(2); // odd-even sorting : + let sort_steps = self.cache_size.next_power_of_two().ilog2(); // Each element is sorted + let shader_invocations = self.cache_size.div_ceil(2); // bitonic sorting : // half as many shaders // per element let workgroup_invocations = shader_invocations.div_ceil(16); - for _ in 0..sort_element_count.div_ceil(2) + for i in 0..sort_steps { - compute_pass.dispatch_workgroups(workgroup_invocations as u32, 1, 1); + for j in 0..=i + { + compute_pass.set_immediates(0, bytes_of(&(j << 16 | i))); + compute_pass.dispatch_workgroups(workgroup_invocations as u32, 1, 1); + } } } } @@ -937,7 +970,7 @@ pub struct VoxelCache // Invalidation invalidation_pipeline: ComputePipeline, - usage_buffer: UsageBuffer, + pub usage_buffer: UsageBuffer, } pub struct StructurePoolElement