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

refactor: reorganising files

Browse Source
This commit is contained in:
octagonal
2026-09-05 16:01:31 +02:00
parent f8cfb69b21
commit a241b7fd83
12 changed files with 2179 additions and 3917 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/main.rs
+8 -36
View File
@@ -1,65 +1,40 @@
#![feature(generic_const_exprs)]
#![feature(float_algebraic)]
use core::sync;
use std::cell::RefCell;
use std::collections::HashMap;
use std::fs::File;
use std::hash::Hash;
use std::rc::Rc;
use std::sync::Arc;
use std::sync::atomic::AtomicBool;
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;
use egui::Label;
use egui::emath::fast_midpoint;
use egui::mutex::Mutex;
use egui_plot::BarChart;
use glam::Mat4;
use glam::Vec3;
use glam::Vec4;
use itertools::Itertools;
use rand::random;
use rayon::iter::IndexedParallelIterator;
use rayon::iter::IntoParallelRefIterator;
use rayon::iter::ParallelIterator;
use wgpu::BindGroup;
use wgpu::BindGroupEntry;
use wgpu::BindGroupLayoutDescriptor;
use wgpu::BindGroupLayoutEntry;
use wgpu::Buffer;
use wgpu::BufferUsages;
use wgpu::DepthBiasState;
use wgpu::Device;
use wgpu::Extent3d;
use wgpu::Features;
use wgpu::FragmentState;
use wgpu::InstanceDescriptor;
use wgpu::InstanceFlags;
use wgpu::MemoryBudgetThresholds;
use wgpu::NoopBackendOptions;
use wgpu::Operations;
use wgpu::PrimitiveState;
use wgpu::RenderPassDepthStencilAttachment;
use wgpu::RenderPipeline;
use wgpu::RenderPipelineDescriptor;
use wgpu::ShaderModuleDescriptor;
use wgpu::ShaderStages;
use wgpu::StencilState;
use wgpu::Texture;
use wgpu::TextureUsages;
use wgpu::TextureView;
use wgpu::VertexState;
use wgpu::util::BufferInitDescriptor;
use wgpu::util::DeviceExt;
use wgpu::util::DownloadBuffer;
use wgpu::util::StagingBelt;
use winit::application::ApplicationHandler;
use winit::event::DeviceEvent;
use winit::event::MouseScrollDelta;
@@ -69,30 +44,27 @@ use winit::event_loop::ActiveEventLoop;
use winit::event_loop::ControlFlow;
use winit::event_loop::EventLoop;
use winit::event_loop::OwnedDisplayHandle;
use winit::platform::x11::EventLoopBuilderExtX11;
use winit::window::Window;
use winit::window::WindowId;
use crate::camera::Camera;
use crate::egui_renderer::EguiRenderer;
use crate::producers::BallGenerator;
use crate::producers::ChunkedProducer;
use crate::producers::Producer;
use crate::producers::SineGenerator;
use crate::producers::TerrainGenerator;
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;
use crate::sparse_tree::NTreeNodeLocator;
use crate::voxel_cache::VoxelCache;
use crate::voxel_cache::data::CacheNodeRequest;
use crate::voxel_cache::data::CacheResponse;
use crate::voxel_cache::data::ColorBytes;
use crate::voxel_cache::data::DestinationElement;
use crate::voxel_cache::data::LocationPoolElement;
mod camera;
mod egui_renderer;
mod producers;
mod sparse_tree;
mod voxel_cache;
//mod tree;
//
struct State
src/producers.rs
+3 -4
View File
@@ -2,12 +2,11 @@ use std::fs::File;
use std::path::Path;
use glam::Vec3;
use indicatif::ProgressIterator;
use itertools::Itertools;
use crate::voxel_cache::gpu::ExplicitNTreeNode;
use crate::voxel_cache::gpu::StructurePointer;
use crate::voxel_cache::sparse::Color;
use crate::sparse_tree::Color;
use crate::voxel_cache::data::ExplicitNTreeNode;
use crate::voxel_cache::data::StructurePointer;
pub struct BallGenerator<const N: usize>
{
src/voxel_cache/sparse.rs → src/sparse_tree.rs
+2 -2
View File
@@ -5,8 +5,8 @@ use bytemuck::Pod;
use bytemuck::Zeroable;
use itertools::Itertools;
use crate::voxel_cache::gpu::ExplicitNTreeNode;
use crate::voxel_cache::gpu::StructurePointer;
use crate::voxel_cache::data::ExplicitNTreeNode;
use crate::voxel_cache::data::StructurePointer;
#[derive(Debug, Clone, Copy, PartialEq, Pod, Zeroable)]
#[repr(C)]
src/voxel_cache.rs
+1042 -4
View File
File diff suppressed because it is too large Load Diff
src/voxel_cache/cache.rs
-2113
View File
File diff suppressed because it is too large Load Diff
src/voxel_cache/cache.wait.rs
-1663
View File
File diff suppressed because it is too large Load Diff
src/voxel_cache/data.rs
+102
View File
@@ -0,0 +1,102 @@
use bytemuck::Pod;
use bytemuck::Zeroable;
use wgpu::Buffer;
use crate::sparse_tree::Color;
#[derive(Clone, Copy, Zeroable, Pod, Debug)]
#[repr(C)]
pub struct CacheNodeRequest
{
// Requested ressource
pub structure_id: u32,
pub structure_locator: u32,
// Write back info
pub node_index: u32,
pub child_index: u32,
}
pub struct RequestBufferElement<const N: usize>
where
[(); N * N * N]:,
{
request_count: [u32; N * N * N],
}
pub struct StructurePoolElement<const N: usize>
where
[(); N * N * N]:,
{
pointers: [StructurePointer; N * N * N],
}
pub struct DestinationElement
{
pub node: u32,
pub child: u32,
}
pub struct ColorBytes(pub u8, pub u8, pub u8, pub u8);
impl From<Color> for ColorBytes
{
fn from(value: Color) -> Self
{
Self(
(value.0 * 255.) as u8,
(value.1 * 255.) as u8,
(value.2 * 255.) as u8,
(value.3 * 255.) as u8,
)
}
}
pub struct ColorPoolElement<const N: usize>
where
[(); N * N * N]:,
{
colors: [ColorBytes; N * N * N],
}
pub struct LocationPoolElement
{
pub structure_id: u32,
pub structure_locator: u32,
}
pub struct CacheResponse
{
// Each buffer contains the same amount of elements (structure of arrays style)
// Cache data to bring in
pub structure_nodes: Buffer,
pub color_nodes: Buffer,
pub locations: Buffer,
// Which nodes this extends : node_index + child_index
pub parents: Buffer,
}
#[derive(Clone, Copy, Pod, Zeroable)]
#[repr(transparent)]
pub struct StructurePointer(pub u32);
impl StructurePointer
{
pub fn new(subdivided: bool, pointer_valid: bool, pointer: u32) -> Self
{
assert!(pointer >> 30 == 0);
StructurePointer((subdivided as u32) << 31 | (pointer_valid as u32) << 30 | pointer)
}
}
#[derive(Clone, Copy, Zeroable)]
#[repr(C)]
pub struct ExplicitNTreeNode<const N: usize>
where
[(); N * N * N]:,
{
pub structure: [StructurePointer; N * N * N],
pub colors: [Color; N * N * N],
}
src/voxel_cache/gpu.rs
-27
View File
@@ -1,27 +0,0 @@
use bytemuck::Pod;
use bytemuck::Zeroable;
use crate::voxel_cache::sparse::Color;
#[derive(Clone, Copy, Pod, Zeroable)]
#[repr(transparent)]
pub struct StructurePointer(pub u32);
impl StructurePointer
{
pub fn new(subdivided: bool, pointer_valid: bool, pointer: u32) -> Self
{
assert!(pointer >> 30 == 0);
StructurePointer((subdivided as u32) << 31 | (pointer_valid as u32) << 30 | pointer)
}
}
#[derive(Clone, Copy, Zeroable)]
#[repr(C)]
pub struct ExplicitNTreeNode<const N: usize>
where
[(); N * N * N]:,
{
pub structure: [StructurePointer; N * N * N],
pub colors: [Color; N * N * N],
}
src/voxel_cache/pipeline.rs
-68
View File
@@ -1,68 +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::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,
}
src/voxel_cache/request_buffer.rs
+644
View File
@@ -0,0 +1,644 @@
use bytemuck::bytes_of;
use wgpu::BindGroup;
use wgpu::BindGroupEntry;
use wgpu::Buffer;
use wgpu::BufferUsages;
use wgpu::CommandEncoder;
use wgpu::ComputePipeline;
use wgpu::Device;
use wgpu::ShaderStages;
use crate::voxel_cache::data::RequestBufferElement;
pub struct RequestBuffer<const N: usize>
{
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,
pub bindgroup: BindGroup,
pub device: Device,
pub compaction_pipeline: ComputePipeline,
pub running_sum_pipeline: ComputePipeline,
}
impl<const N: usize> RequestBuffer<N>
where
[(); N * N * N]:,
{
pub fn new(cache_size: usize, device: Device) -> Self
{
let request_buffer = device.create_buffer(&wgpu::BufferDescriptor {
label: Some(format!("request_buffer_{N}").as_str()),
size: (size_of::<RequestBufferElement<N>>() * cache_size) as u64,
usage: BufferUsages::STORAGE | BufferUsages::COPY_DST,
mapped_at_creation: false,
});
let request_count_buffer = device.create_buffer(&wgpu::BufferDescriptor {
label: Some("Request_count_buffer"),
size: size_of::<u32>() as u64,
usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
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::<wgpu::util::DispatchIndirectArgs>() 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::<wgpu::util::DispatchIndirectArgs>() 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::<u32>() 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;
let sort_buffer = device.create_buffer(&wgpu::BufferDescriptor {
label: Some(format!("request_sort_buffer_{N}").as_str()),
size: (size_of::<(u32, u32)>() * sort_buffer_count) as u64,
usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
mapped_at_creation: true,
});
let init_data = (0..(cache_size as u32))
.flat_map(|i| (0..((N * N * N) as u32)).map(move |j| (i, j)))
.flat_map(|(i, j)| [i, j]) // Because bytemuck does not like tuples ...
.collect::<Vec<_>>();
sort_buffer
.get_mapped_range_mut(0..)
.unwrap()
.copy_from_slice(bytemuck::cast_slice(&init_data));
sort_buffer.unmap();
let bindgroup_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
label: Some("request_sort_pipeline_bing_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,
},
wgpu::BindGroupLayoutEntry {
binding: 1,
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: 2,
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: 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,
},
],
});
let bindgroup = device.create_bind_group(&wgpu::BindGroupDescriptor {
label: Some("request_sort_pipeline_bindgroup"),
layout: &bindgroup_layout,
entries: &[
BindGroupEntry {
binding: 0,
resource: request_buffer.as_entire_binding(),
},
BindGroupEntry {
binding: 1,
resource: sort_buffer.as_entire_binding(),
},
BindGroupEntry {
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(),
},
],
});
let pipeline_layouts = &device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
label: Some("request_buffer_pipeline_layout"),
bind_group_layouts: &[Some(&bindgroup_layout)],
immediate_size: size_of::<u32>() as u32,
});
let children_count = N * N * N;
let sort_pipeline =
device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
label: Some("Sort node request"),
layout: Some(pipeline_layouts),
module: &device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: Some("clear_chunk_requests shader module"),
source: wgpu::ShaderSource::Wgsl(
format!("
struct RequestElement
{{
children: array<u32, {children_count}>
}}
struct SortElement
{{
node: u32,
child: u32
}}
var<immediate> both_phase: u32;
@group(0) @binding(0) var<storage, read_write> request_buffer: array<RequestElement>;
@group(0) @binding(1) var<storage, read_write> sort_indirection: array<SortElement>;
@group(0) @binding(2) var<storage, read_write> request_count: u32;
@group(0) @binding(3) var<storage, read_write> indirect_count: vec3<u32>;
@group(0) @binding(4) var<storage, read_write> element_sort_count: 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(64)
fn main(
@builtin(global_invocation_id) global_invocation_id: vec3<u32>
)
{{
let index = global_invocation_id.x;
let total = arrayLength(&sort_indirection);
let sub_phase = (both_phase >> 16) & 0xFFFF;
let phase = both_phase & 0xFFFF;
let phase_total_width = u32(1 << (phase + 1));
// Check if phase is last
/*
if(phase == sub_phase)
{{
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<u32>(0);
}}
}}
*/
if(phase == 0)
{{
let a = index * 2;
let b = a + 1;
sort_indirection[a].child = a % {children_count};
sort_indirection[b].child = b % {children_count};
}}
var swap_indices = vec2<u32>(0, 0);
if(sub_phase == 0)
{{
// Bitonic fusion
swap_indices = big_fusion(index, phase_total_width);
}}else
{{
swap_indices = small_fusion(index, phase, sub_phase);
}}
// Do swap
//if(swap_indices.y >= request_count_round_up)
if(swap_indices.y >= element_sort_count)
{{
// Suppose that swap_indices.y is -inf, dont swap
return;
}}
let av = request_buffer[sort_indirection[swap_indices.x].node].children[sort_indirection[swap_indices.x].child];
let bv = request_buffer[sort_indirection[swap_indices.y].node].children[sort_indirection[swap_indices.y].child];
if(bv > av)
{{
let temp = sort_indirection[swap_indices.x];
sort_indirection[swap_indices.x] = sort_indirection[swap_indices.y];
sort_indirection[swap_indices.y] = temp;
}}
}}
")
.into(),
),
}),
entry_point: Some("main"),
compilation_options: Default::default(),
cache: None,
});
let running_sum_pipeline =
device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
label: Some("running_sum_pipeline"),
layout: Some(pipeline_layouts),
module: &device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: Some("running_sum shader module"),
source: wgpu::ShaderSource::Wgsl(
format!("
struct RequestElement
{{
children: array<u32, {children_count}>
}}
struct SortElement
{{
node: u32,
child: u32
}}
var<immediate> phase: u32;
@group(0) @binding(0) var<storage, read_write> request_buffer: array<RequestElement>;
@group(0) @binding(1) var<storage, read_write> sort_indirection: array<SortElement>;
@group(0) @binding(2) var<storage, read_write> request_counts: atomic<u32>;
@group(0) @binding(3) var<storage, read_write> indirect_count: u32;
@group(0) @binding(4) var<storage, read_write> node_sort_count: u32;
@compute
@workgroup_size(64)
fn main(
@builtin(global_invocation_id) global_invocation_id: vec3<u32>
)
{{
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 == 0xFFFFFFFF)
{{
// Double buffering bring back
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(),
),
}),
entry_point: Some("main"),
compilation_options: Default::default(),
cache: None,
});
let compaction_pipeline =
device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
label: Some("compaction_pipeline"),
layout: Some(pipeline_layouts),
module: &device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: Some("compaction shader module"),
source: wgpu::ShaderSource::Wgsl(
format!("
struct RequestElement
{{
children: array<u32, {children_count}>
}}
struct SortElement
{{
node: u32,
child: u32
}}
var<immediate> phase: u32;
@group(0) @binding(0) var<storage, read_write> request_buffer: array<RequestElement>;
@group(0) @binding(1) var<storage, read_write> sort_indirection: array<SortElement>;
@group(0) @binding(2) var<storage, read_write> request_counts: atomic<u32>;
@group(0) @binding(3) var<storage, read_write> indirect_count: vec3<u32>;
@group(0) @binding(4) var<storage, read_write> element_sort_count: u32;
@compute
@workgroup_size(64)
fn main(
@builtin(global_invocation_id) global_invocation_id: vec3<u32>
)
{{
let index = global_invocation_id.x;
let sindex = index * {children_count};
let len = arrayLength(&request_buffer);
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++)
{{
count += request_buffer[index].children[i];
}}
if(count != 0) // Keep ?
{{
for(var i = u32(0); i < u32({children_count}); i++)
{{
sort_indirection[destination_index * u32({children_count}) + i].node = index;
}}
}}
}}
")
.into(),
),
}),
entry_point: Some("main"),
compilation_options: Default::default(),
cache: None,
});
let reset_pipeline =
device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
label: Some("reset_node_requests"),
layout: Some(pipeline_layouts),
module: &device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: Some("clear_chunk_requests shader module"),
source: wgpu::ShaderSource::Wgsl(
format!("
struct RequestElement
{{
children: array<u32, {children_count}>
}}
@group(0) @binding(0) var<storage, read_write> request_buffer: array<RequestElement>;
@group(0) @binding(1) var<storage, read_write> _ignore: array<u32>;
@group(0) @binding(2) var<storage, read_write> request_count: u32;
@group(0) @binding(4) var<storage, read_write> node_sort_count: u32;
@compute
@workgroup_size(64)
fn main(
@builtin(global_invocation_id) global_invocation_id: vec3<u32>
)
{{
let index = global_invocation_id.x;
let total = arrayLength(&request_buffer);
if(index >= total)
{{ return; }}
if(index < total)
{{
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;
}}
}}
")
.into(),
),
}),
entry_point: Some("main"),
compilation_options: Default::default(),
cache: None,
});
Self {
cache_size,
request_buffer,
request_count_buffer,
indirect_count,
indirect_count_storage,
element_sort_count_buffer,
running_sum_pipeline,
sort_buffer,
sort_pipeline,
reset_pipeline,
compaction_pipeline,
bindgroup,
device,
}
}
pub fn request_buffer(&self) -> &Buffer
{
&self.request_buffer
}
pub fn request_count_buffer(&self) -> &Buffer
{
&self.request_count_buffer
}
pub fn sort_buffer(&self) -> &Buffer
{
&self.sort_buffer
}
pub fn reset_requests(&self, encoder: &mut CommandEncoder)
{
let mut compute_pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
label: Some("request_buffer_reset_compute_pass"),
timestamp_writes: None,
});
compute_pass.set_bind_group(0, Some(&self.bindgroup), &[]);
compute_pass.set_pipeline(&self.reset_pipeline);
let shader_invocations = self.cache_size; // one invocation per element
let workgroup_invocations = shader_invocations.div_ceil(64);
compute_pass.dispatch_workgroups(workgroup_invocations as u32, 1, 1);
}
pub fn sort_requests(&self, encoder: &mut CommandEncoder)
{
let mut compaction_compute_pass =
encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
label: Some("request_buffer_compaction_compute_pass"),
timestamp_writes: None,
});
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
compaction_compute_pass.set_pipeline(&self.running_sum_pipeline);
// Phase 0: Put ones in correct location
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
let running_sum_steps = request_element_count.next_power_of_two().ilog2();
for i in 1..=running_sum_steps
{
compaction_compute_pass.set_immediates(0, bytes_of(&i));
compaction_compute_pass.dispatch_workgroups(workgroups_invocations as u32, 1, 1);
}
if !running_sum_steps.is_multiple_of(2)
{
// 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
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
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::<wgpu::util::DispatchIndirectArgs>() 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
{
sort_compute_pass.set_immediates(0, bytemuck::bytes_of(&(i | (j << 16))));
sort_compute_pass.dispatch_workgroups_indirect(&self.indirect_count, 0);
}
drop(sort_compute_pass);
}
}
}
src/voxel_cache/structure_table.rs
+129
View File
@@ -0,0 +1,129 @@
use std::num::NonZero;
use wgpu::Buffer;
use wgpu::BufferUsages;
use wgpu::CommandEncoder;
use wgpu::Device;
use wgpu::Queue;
use wgpu::util::StagingBelt;
use crate::voxel_cache::data::StructurePointer;
use crate::voxel_cache::request_buffer::RequestBuffer;
// Stores root pointers to the cache
pub struct StructureTable
{
pub(crate) device: Device,
pub(crate) queue: Queue,
pub(crate) allocation_table: Vec<bool>,
pub(crate) available_slots: usize,
pub(crate) pointer_table: Buffer,
pub(crate) request_buffer: RequestBuffer<1>,
pub(crate) write_staging: StagingBelt,
}
impl StructureTable
{
pub fn new(device: Device, queue: Queue) -> Self
{
let pointer_table = device.create_buffer(&wgpu::BufferDescriptor {
label: Some("structure_table_pointer_table"),
size: size_of::<u32>() as u64,
usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC,
mapped_at_creation: false,
});
StructureTable {
request_buffer: RequestBuffer::new(1, device.clone()),
write_staging: StagingBelt::new(device.clone(), size_of::<u32>() as u64),
device,
queue,
allocation_table: vec![false],
available_slots: 1,
pointer_table,
}
}
fn double_capacity(&mut self, encoder: &mut CommandEncoder)
{
self.available_slots += self.allocation_table.len();
self.allocation_table
.extend(vec![false; self.allocation_table.len()]);
// Copy buffers into bigger buffers
let pointer_table = self.device.create_buffer(&wgpu::BufferDescriptor {
label: Some("structure_table_pointer_table"),
size: size_of::<u32>() as u64 * self.allocation_table.len() as u64,
usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC | BufferUsages::COPY_DST,
mapped_at_creation: false,
});
encoder.copy_buffer_to_buffer(
&self.pointer_table,
0,
&pointer_table,
0,
self.pointer_table.size(),
);
self.pointer_table = pointer_table;
self.request_buffer = RequestBuffer::new(self.allocation_table.len(), self.device.clone());
}
pub fn request_buffer(&self) -> &RequestBuffer<1>
{
&self.request_buffer
}
pub fn request_buffer_mut(&mut self) -> &mut RequestBuffer<1>
{
&mut self.request_buffer
}
pub fn remove_structure(&mut self, structure_id: u32)
{
self.available_slots += 1;
self.allocation_table[structure_id as usize] = false;
}
// Returns a structure Id
pub fn allocate_structure(&mut self, subdivided: bool) -> u32
{
let mut encoder = self
.device
.create_command_encoder(&wgpu::CommandEncoderDescriptor {
label: Some("structure_table_resize_encoder"),
});
// Get first available element
if self.available_slots == 0
{
self.double_capacity(&mut encoder);
}
let (first_id, _) = self
.allocation_table
.iter()
.enumerate()
.filter(|(_, allocated)| !**allocated)
.next()
.unwrap();
self.allocation_table[first_id] = true;
self.available_slots -= 1;
// Write empty pointer to new pointer
self.write_staging
.write_buffer(
&mut encoder,
&self.pointer_table,
size_of::<u32>() as u64 * first_id as u64,
NonZero::new(size_of::<u32>() as u64).unwrap(),
)
.copy_from_slice(bytemuck::bytes_of(
&StructurePointer::new(subdivided, false, 0).0,
));
self.write_staging.finish_and_recall_on_submit(&encoder);
self.queue.submit([encoder.finish()]);
first_id as u32
}
}
src/voxel_cache/usage_buffer.rs
+249
View File
@@ -0,0 +1,249 @@
use bytemuck::bytes_of;
use wgpu::BindGroup;
use wgpu::BindGroupEntry;
use wgpu::Buffer;
use wgpu::BufferUsages;
use wgpu::CommandEncoder;
use wgpu::ComputePipeline;
use wgpu::Device;
use wgpu::ShaderStages;
pub struct UsageBuffer
{
pub cache_size: usize,
pub current_timestamp: u32,
pub usage_buffer: Buffer,
pub sort_buffer: Buffer,
pub sort_pipeline: ComputePipeline,
pub bindgroup: BindGroup,
pub device: Device,
}
impl UsageBuffer
{
pub fn new(cache_size: usize, device: Device) -> Self
{
let initial_timestamp = 0;
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 | BufferUsages::COPY_SRC,
mapped_at_creation: true,
});
let sort_buffer = device.create_buffer(&wgpu::BufferDescriptor {
label: Some("usage_sort_buffer"),
size: (size_of::<u32>() * cache_size) as u64,
usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
mapped_at_creation: true,
});
usage_buffer
.get_mapped_range_mut(0..)
.unwrap()
.copy_from_slice(bytemuck::cast_slice(
vec![initial_timestamp; cache_size].as_slice(),
));
usage_buffer.unmap();
sort_buffer
.get_mapped_range_mut(0..)
.unwrap()
.copy_from_slice(bytemuck::cast_slice(
&(0..(cache_size as u32)).collect::<Vec<_>>(),
));
sort_buffer.unmap();
let bindgroup_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
label: Some("usage_sort_pipeline_bing_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,
},
wgpu::BindGroupLayoutEntry {
binding: 1,
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 bindgroup = device.create_bind_group(&wgpu::BindGroupDescriptor {
label: Some("usage_sort_pipeline_bindgroup"),
layout: &bindgroup_layout,
entries: &[
BindGroupEntry {
binding: 0,
resource: usage_buffer.as_entire_binding(),
},
BindGroupEntry {
binding: 1,
resource: sort_buffer.as_entire_binding(),
},
],
});
let pipeline_layouts = &device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
label: Some("usage_buffer_pipeline_layout"),
bind_group_layouts: &[Some(&bindgroup_layout)],
immediate_size: size_of::<u32>() as u32,
});
let sort_pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
label: Some("Sort node request"),
layout: Some(pipeline_layouts),
module: &device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: Some("clear_chunk_requests shader module"),
source: wgpu::ShaderSource::Wgsl(
"
@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(64)
fn main(
@builtin(global_invocation_id) global_invocation_id: vec3<u32>
)
{{
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));
var swap_indices = vec2<u32>(0, 0);
if(sub_phase == 0)
{{
// Bitonic fusion
swap_indices = big_fusion(index, phase_total_width);
}}else
{{
swap_indices = small_fusion(index, phase, sub_phase);
}}
// 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]];
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(),
),
}),
entry_point: Some("main"),
compilation_options: Default::default(),
cache: None,
});
Self {
cache_size,
current_timestamp: initial_timestamp,
usage_buffer,
sort_buffer,
sort_pipeline,
bindgroup,
device,
}
}
pub fn timestamp(&self) -> u32
{
self.current_timestamp
}
pub fn next_frame(&mut self)
{
let (new_timestamp, _) = self.current_timestamp.overflowing_add(1);
self.current_timestamp = new_timestamp;
}
pub fn usage_buffer(&self) -> &Buffer
{
&self.usage_buffer
}
pub fn sort_buffer(&self) -> &Buffer
{
&self.sort_buffer
}
pub fn sort_usage(&self, encoder: &mut CommandEncoder)
{
let mut compute_pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
label: Some("usage_buffer_sort_compute_pass"),
timestamp_writes: None,
});
compute_pass.set_bind_group(0, Some(&self.bindgroup), &[]);
compute_pass.set_pipeline(&self.sort_pipeline);
// Compute required shader invocations
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(64);
for i in 0..=sort_steps
{
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);
}
}
}
}