octagonal/vxls
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Adding histogram

Browse Source
This commit is contained in:
octagonal
2026-09-02 15:39:25 +02:00
parent 194cbf04ef
commit 2fbce4ab81
3 changed files with 118 additions and 44 deletions
Download Patch File Download Diff File Expand all files Collapse all files
Cargo.toml
+1
View File
@@ -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"
src/main.rs
+47 -7
View File
@@ -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<Mutex<VoxelCache<4>>>,
usage_vec: Arc<Mutex<Vec<usize>>>,
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);
}
}
src/voxel/cache.rs
+70 -37
View File
@@ -601,7 +601,7 @@ impl UsageBuffer
let usage_buffer = device.create_buffer(&wgpu::BufferDescriptor {
label: Some("Usage buffer"),
size: (size_of::<u32>() * 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::<u32>() 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<storage, read_write> usage_stamps: array<u32>;
@group(0) @binding(1) var<storage, read_write> sort_indirection: array<u32>;
@group(0) @binding(0) var<storage, read_write> usage_buffer: array<u32>;
@group(0) @binding(1) var<storage, read_write> sort_buffer: array<u32>;
var<immediate> both_phase: u32;
fn big_fusion(index: u32, phase_size: u32) -> vec2<u32>
{{
// 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<u32>(a, b);
}}
fn small_fusion(index: u32, phase: u32, sub_phase: u32) -> vec2<u32>
{{
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<u32>(a, b);
}}
@compute
@workgroup_size(16)
@@ -691,39 +719,40 @@ impl UsageBuffer
@builtin(global_invocation_id) global_invocation_id: vec3<u32>
)
{{
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<u32>(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<const N: usize>
// Invalidation
invalidation_pipeline: ComputePipeline,
usage_buffer: UsageBuffer,
pub usage_buffer: UsageBuffer,
}
pub struct StructurePoolElement<const N: usize>