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

Works

Browse Source
This commit is contained in:
octagonal
2026-09-05 15:39:19 +02:00
parent 7efb1a1404
commit f8cfb69b21
14 changed files with 495 additions and 960 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/main.rs
+68 -20
View File
@@ -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<Mutex<Vec<usize>>>,
rm_time: Arc<Mutex<f32>>,
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::<u64>() * 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: &timestamp_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(&timestamp_query, 0..2, &timestamp_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,
&timestamp_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
src/producers.rs
+166 -25
View File
@@ -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<const N: usize>
{
@@ -15,7 +16,13 @@ pub struct BallGenerator<const N: usize>
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<const N: usize>
@@ -240,6 +247,14 @@ where
heightmap: Vec<f32>,
colormap: Vec<u8>,
heightmap_low_width: usize,
heightmap_low_height: usize,
heightmap_low: Vec<(f32, f32)>,
colormap_low_width: usize,
colormap_low_height: usize,
colormap_mip: Vec<u8>,
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;
src/voxel/pipeline.rs
-799
View File
@@ -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<const N: usize>
where
[(); N * N * N]:,
{
chunk_allocations: Vec<bool>,
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<const N: usize> VoxelPipeline<N>
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::<u32>() 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::<u32>() 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::<CacheChunkObject>() 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::<u32>() 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::<StructurePoolElement<N>>() 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::<ColorPoolElement<N>>() 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::<LocationPoolElement>() 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::<RequestBufferElement<N>>() 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::<u32>() 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::<Vec<_>>().as_slice(),
));
request_sort_buffer.unmap();
let usage_buffer = device.create_buffer(&wgpu::wgt::BufferDescriptor {
label: Some("Usage buffer buffer"),
size: size_of::<u32>() 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<storage, read_write> chunk_requests: array<u32>;
@compute
@workgroup_size(16)
fn main(
@builtin(global_invocation_id) global_invocation_id: vec3<u32>
)
{
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<atomic<u32>, 64>
}
@group(0) @binding(0) var<storage, read_write> structure_pool: array<u32>;
@group(0) @binding(1) var<storage, read_write> color_pool: array<u32>;
@group(0) @binding(2) var<storage, read_write> location_pool: array<u32>;
@group(0) @binding(3) var<storage, read_write> request_buffer: array<RequestElement>;
@group(0) @binding(4) var<storage, read_write> usage_buffer: array<u32>;
@group(1) @binding(0) var<storage, read_write> sort_indirection: array<u32>;
@compute
@workgroup_size(16)
fn main(
@builtin(global_invocation_id) global_invocation_id: vec3<u32>
)
{
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::<CacheChunkObject>() 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::<Vec<_>>();
self.chunk_indices = self.device.create_buffer(&wgpu::wgt::BufferDescriptor {
label: Some("Chunk index buffer"),
size: size_of::<u32>() 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::<u32>() 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<ChunkHandle>
{
// 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::<ChunkObject>() 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::<u32>() 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::<CacheChunkObject>() as u64,
NonZero::new(size_of::<CacheChunkObject>() 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()
}
}
src/voxel.rs → src/voxel_cache.rs
View File
src/voxel/cache.rs → src/voxel_cache/cache.rs
+182 -107
View File
@@ -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<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,
@@ -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::<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;
@@ -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<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>
{{
@@ -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<u32>(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<u32>(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<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: u32;
@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)
@@ -321,57 +372,55 @@ where
@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);
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<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)
@@ -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<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: atomic<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)
@@ -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::<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
{
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::<CacheNodeRequest>() 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
src/voxel/cache.wait.rs → src/voxel_cache/cache.wait.rs
View File
src/voxel/gpu.rs → src/voxel_cache/gpu.rs
+1 -1
View File
@@ -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)]
src/voxel_cache/pipeline.rs
+68
View File
@@ -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,
}
src/voxel_cache/request_buffer.rs
View File
src/voxel/sparse.rs → src/voxel_cache/sparse.rs
+2 -2
View File
@@ -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)]
src/voxel_cache/structure_table.rs
View File
src/voxel_cache/usage_buffer.rs
View File