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

Terrain !

Browse Source
This commit is contained in:
octagonal
2026-09-03 23:23:43 +02:00
parent 2fbce4ab81
commit b826de15a6
6 changed files with 902 additions and 177 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/camera.rs
+1 -1
View File
@@ -18,7 +18,7 @@ pub struct Camera
fov: f32,
pub aspect: f32,
speed: f32,
pub speed: f32,
pub pressed_keyset: HashSet<KeyCode>,
}
src/main.rs
+197 -62
View File
@@ -4,6 +4,7 @@ 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;
@@ -13,6 +14,8 @@ use bytemuck::Pod;
use bytemuck::Zeroable;
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;
@@ -21,6 +24,9 @@ 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;
@@ -54,6 +60,7 @@ use wgpu::util::DownloadBuffer;
use wgpu::util::StagingBelt;
use winit::application::ApplicationHandler;
use winit::event::DeviceEvent;
use winit::event::MouseScrollDelta;
use winit::event::WindowEvent;
use winit::event_loop;
use winit::event_loop::ActiveEventLoop;
@@ -66,9 +73,14 @@ 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::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;
@@ -103,6 +115,10 @@ struct State
pipeline: RenderPipeline,
voxel_cache: Arc<Mutex<VoxelCache<4>>>,
terrain_generator: Arc<TerrainGenerator<4>>,
chunk_pos_map: Arc<HashMap<u32, (usize, usize, usize)>>,
instance_buffer: Buffer,
instance_count: usize,
usage_vec: Arc<Mutex<Vec<usize>>>,
insertion_debounce: bool,
@@ -123,6 +139,16 @@ struct Immediates
frame_timestamp: u32,
}
#[derive(Debug, Clone, Copy, Zeroable, Pod)]
#[repr(C)]
struct InstanceAttribute
{
x: f32,
y: f32,
z: f32,
id: u32,
}
impl State
{
async fn new(display: OwnedDisplayHandle, window: Arc<Window>) -> State
@@ -141,7 +167,7 @@ impl State
.unwrap();
let (device, queue) = adapter
.request_device(&wgpu::DeviceDescriptor {
required_features: Features::IMMEDIATES,
required_features: Features::IMMEDIATES | Features::SHADER_EARLY_DEPTH_TEST,
required_limits: wgpu::Limits {
max_immediate_size: 112,
max_storage_buffers_per_shader_stage: 16,
@@ -160,9 +186,34 @@ impl State
let egui_renderer = EguiRenderer::new(&device, surface_format, &window);
let mut voxel_cache = VoxelCache::<4>::new(16000, device.clone(), queue.clone());
let id = voxel_cache.structure_table.allocate_structure(true);
println!("id: {id}");
let mut voxel_cache = VoxelCache::<4>::new(100_000, device.clone(), queue.clone());
let terrain_generator = TerrainGenerator::<4>::new(5, "vxls_height.tif", 0.2, "img.jpg");
let mut chunk_pos_map = HashMap::new();
let chunk_instances = (0..terrain_generator.chunk_width)
.cartesian_product(0..terrain_generator.chunk_height)
.cartesian_product(0..terrain_generator.chunk_alt)
.map(|((x, z), y)| {
let chunk_id = voxel_cache.structure_table.allocate_structure(true);
//dbg!(chunk_id);
chunk_pos_map.insert(chunk_id, (x, y, z));
InstanceAttribute {
x: x as f32,
y: y as f32,
z: z as f32,
id: chunk_id,
}
})
.collect::<Vec<_>>();
let instance_count = chunk_instances.len();
let instance_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
label: Some("Instance buffer"),
contents: bytemuck::cast_slice(chunk_instances.as_slice()),
usage: BufferUsages::COPY_DST | BufferUsages::VERTEX,
});
let shader_module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: Some("Main shader module"),
@@ -187,13 +238,28 @@ impl State
module: &shader_module,
entry_point: Some("chunk"),
compilation_options: Default::default(),
buffers: &[],
buffers: &[Some(wgpu::VertexBufferLayout {
array_stride: (size_of::<f32>() * 3 + size_of::<u32>()) as u64,
step_mode: wgpu::VertexStepMode::Instance,
attributes: &[
wgpu::VertexAttribute {
format: wgpu::VertexFormat::Float32x3,
offset: 0,
shader_location: 0,
},
wgpu::VertexAttribute {
format: wgpu::VertexFormat::Uint32,
offset: (size_of::<f32>() * 3) as u64,
shader_location: 1,
},
],
})],
},
primitive: wgpu::PrimitiveState {
topology: wgpu::PrimitiveTopology::TriangleList,
strip_index_format: None,
front_face: wgpu::FrontFace::Ccw,
cull_mode: None,
cull_mode: Some(wgpu::Face::Front),
unclipped_depth: false,
polygon_mode: wgpu::PolygonMode::Fill,
conservative: false,
@@ -201,7 +267,7 @@ impl State
depth_stencil: Some(wgpu::DepthStencilState {
format: wgpu::TextureFormat::Depth24PlusStencil8,
depth_write_enabled: Some(true),
depth_compare: Some(wgpu::CompareFunction::LessEqual),
depth_compare: Some(wgpu::CompareFunction::Less),
stencil: wgpu::StencilState::default(),
bias: wgpu::DepthBiasState::default(),
}),
@@ -235,6 +301,10 @@ impl State
voxel_cache: Arc::new(Mutex::new(voxel_cache)),
insertion_debounce: false,
camera: Default::default(),
instance_buffer,
instance_count,
terrain_generator: Arc::new(terrain_generator),
chunk_pos_map: chunk_pos_map.into(),
};
// Configure surface for the first time
@@ -323,9 +393,7 @@ impl State
fn render(&mut self)
{
self.camera.update();
// Write random shit in request buffer
let mut belt = StagingBelt::new(self.device.clone(), size_of::<u32>() as u64);
self.voxel_cache.lock().next_frame();
// Create texture view.
// NOTE: We must handle Timeout because the surface may be unavailable
@@ -395,6 +463,7 @@ impl State
});
renderpass.set_pipeline(&self.pipeline);
renderpass.set_vertex_buffer(0, self.instance_buffer.slice(..));
renderpass.set_bind_group(0, Some(&self.voxel_cache.lock().bind_group()), &[]);
let imm = [Immediates {
view_proj: self.camera.view_proj(),
@@ -402,7 +471,7 @@ impl State
frame_timestamp: self.voxel_cache.lock().current_timestamp(),
}];
renderpass.set_immediates(0, unsafe { as_raw_bytes(&imm) });
renderpass.draw(0..36, 0..1);
renderpass.draw(0..36, 0..(self.instance_count as u32));
// End the renderpass.
drop(renderpass);
@@ -410,14 +479,20 @@ impl State
let requests = self.device.create_buffer(&wgpu::BufferDescriptor {
label: Some("dummy_dumb_dinky_aaaahhh_buffer"),
size: 16 * 256,
size: 16 * 1024,
usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
mapped_at_creation: false,
});
self.voxel_cache
.lock()
.cache_post_render(&mut encoder, requests.clone());
if !self
.camera
.pressed_keyset
.contains(&winit::keyboard::KeyCode::KeyF)
{
self.voxel_cache
.lock()
.cache_post_render(&mut encoder, requests.clone());
}
// If you wanted to call any drawing commands, they would go here.
{
@@ -425,6 +500,17 @@ impl State
egui::Window::new("Window ! ").resizable(true).show(
self.egui_renderer.context(),
|ui| {
if self
.camera
.pressed_keyset
.contains(&winit::keyboard::KeyCode::KeyF)
{
ui.label(
egui::RichText::new("Cache paused")
.color(Color32::RED)
.size(28.),
);
}
egui_plot::Plot::new("Plot").show(ui, |plot_ui| {
plot_ui.bar_chart(BarChart::new(
"histo",
@@ -488,17 +574,24 @@ impl State
self.insertion_debounce = false;
}
if (self
// if (self
// .camera
// .pressed_keyset
// .contains(&winit::keyboard::KeyCode::KeyF))
// && !self.insertion_debounce
// {
if !self
.camera
.pressed_keyset
.contains(&winit::keyboard::KeyCode::KeyF))
&& !self.insertion_debounce
.contains(&winit::keyboard::KeyCode::KeyF)
{
self.insertion_debounce = true;
let request_count = self.voxel_cache.lock().total_request_count();
let cloned_cache = self.voxel_cache.clone();
let cloned_device = self.device.clone();
let cloned_queue = self.queue.clone();
let cloned_generator = self.terrain_generator.clone();
let cloned_map = self.chunk_pos_map.clone();
let (tx, rx) = sync_channel(1);
DownloadBuffer::read_buffer(
&self.device.clone(),
@@ -513,58 +606,88 @@ impl State
let cache_node_requests: Vec<CacheNodeRequest> =
bytemuck::pod_collect_to_vec(&buffer.unwrap());
let mut sine_gen = SineGenerator::<4>::new(4);
let generator = cloned_generator;
let gen_test = SineGenerator::<4>::new(5);
let mut structure_nodes = vec![];
let mut color_nodes = vec![];
let mut color_nodes: Vec<[ColorBytes; 64]> = vec![];
let mut location_nodes = vec![];
let mut destinations = vec![];
for request in cache_node_requests.iter().take(request_count as usize)
{
let location;
let node;
if request.child_index == u32::MAX
{
// Produce root node
node = sine_gen.produce_node(0, 0, 0, 0);
location = LocationPoolElement {
structure_id: 0,
structure_locator: NTreeNodeLocator::<4>::root().as_usize() as u32,
};
}
else
{
// Figure out depth of request
let locator = NTreeNodeLocator::<4>::from_usize(
request.structure_locator as usize,
);
let depth = locator.depth();
//let (x, y, z) = locator.node_location();
let child_x = request.child_index / (4 * 4);
let child_y = (request.child_index % (4 * 4)) / 4;
let child_z = request.child_index % 4;
cache_node_requests
.par_iter()
.take(request_count as usize)
.map(|request| {
// Unpack structure id
let (chunk_x, chunk_y, chunk_z) =
*cloned_map.get(&request.structure_id).unwrap();
let child_locator =
locator.child(child_x as usize, child_y as usize, child_z as usize);
let location;
let node;
if request.child_index == u32::MAX
{
// Produce root node
node =
generator.produce_node(0, 0, 0, 0, (chunk_x, chunk_y, chunk_z));
location = LocationPoolElement {
structure_id: request.structure_id,
structure_locator: NTreeNodeLocator::<4>::root().as_usize()
as u32,
};
}
else
{
// Figure out depth of request
let locator = NTreeNodeLocator::<4>::from_usize(
request.structure_locator as usize,
);
let depth = locator.depth();
//let (x, y, z) = locator.node_location();
let (nx, ny, nz) = child_locator.node_location();
let child_z = request.child_index / (4 * 4);
let child_y = (request.child_index % (4 * 4)) / 4;
let child_x = request.child_index % 4;
node = sine_gen.produce_node(depth + 1, nx, ny, nz);
let child_locator = locator.child(
child_x as usize,
child_y as usize,
child_z as usize,
);
location = LocationPoolElement {
structure_id: 0,
structure_locator: child_locator.as_usize() as u32,
};
}
let (nx, ny, nz) = child_locator.node_location();
structure_nodes.push(node.structure);
color_nodes.push(node.colors);
location_nodes.push(location);
destinations.push(DestinationElement {
node: request.node_index,
child: request.child_index,
node = generator.produce_node(
depth + 1,
nx,
ny,
nz,
(chunk_x, chunk_y, chunk_z),
);
location = LocationPoolElement {
structure_id: request.structure_id,
structure_locator: child_locator.as_usize() as u32,
};
}
(
node.structure,
node.colors,
location,
DestinationElement {
node: request.node_index,
child: request.child_index,
},
)
})
.collect::<Vec<_>>()
.into_iter()
.for_each(|data| {
structure_nodes.push(data.0);
color_nodes.push(std::array::from_fn(|i| data.1[i].into()));
location_nodes.push(data.2);
destinations.push(data.3);
});
}
if structure_nodes.len() != 0
{
@@ -607,7 +730,7 @@ impl State
}
},
);
println!("Total request count: {}", request_count);
//println!("Total request count: {}", request_count);
loop
{
if rx.try_recv().is_ok()
@@ -617,10 +740,17 @@ impl State
let _ = self.device.poll(wgpu::wgt::PollType::Poll);
}
self.voxel_cache.lock().next_frame();
drop(rx);
}
}
pub fn mouse_wheel(&mut self, delta: MouseScrollDelta)
{
if let MouseScrollDelta::LineDelta(_, y) = delta
{
self.camera.speed += y * (self.camera.speed * 0.05);
}
}
}
#[derive(Default)]
@@ -679,6 +809,11 @@ impl ApplicationHandler for App
// here as this event is always followed up by redraw request.
state.resize(size);
}
WindowEvent::MouseWheel { delta, .. } =>
{
state.mouse_wheel(delta);
}
_ => (),
}
}
src/producers.rs
+377 -16
View File
@@ -1,3 +1,6 @@
use std::fs::File;
use std::path::Path;
use glam::Vec3;
use itertools::Itertools;
@@ -5,7 +8,7 @@ use crate::voxel::gpu::ExplicitNTreeNode;
use crate::voxel::gpu::StructurePointer;
use crate::voxel::sparse::Color;
pub struct SineGenerator<const N: usize>
pub struct BallGenerator<const N: usize>
{
chunk_power: usize,
}
@@ -15,22 +18,28 @@ 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
}
impl<const N: usize> SineGenerator<N>
pub trait Producer<const N: usize>
where
[(); N * N * N]:,
{
fn produce_node(&self, depth: usize, nx: usize, ny: usize, nz: usize) -> ExplicitNTreeNode<N>;
}
impl<const N: usize> BallGenerator<N>
where
[(); N * N * N]:,
{
pub fn new(chunk_power: usize) -> Self
{
SineGenerator { chunk_power }
BallGenerator { chunk_power }
}
}
pub fn produce_node(
&self,
depth: usize,
nx: usize,
ny: usize,
nz: usize,
) -> ExplicitNTreeNode<N>
impl<const N: usize> Producer<N> for BallGenerator<N>
where
[(); N * N * N]:,
{
fn produce_node(&self, depth: usize, nx: usize, ny: usize, nz: usize) -> ExplicitNTreeNode<N>
{
let node_size = N.pow((self.chunk_power - depth) as u32);
let child_size = node_size / N;
@@ -49,23 +58,38 @@ where
let gvy = gny + cy * child_size + (child_size / 2);
let gvz = gnz + cz * child_size + (child_size / 2);
let dist = Vec3::new(gvx as f32 - 128., gvy as f32 - 128., gvz as f32 - 128.).length();
let dist = Vec3::new(
gvx as f32 - (global_size / 2) as f32,
gvy as f32 - (global_size / 2) as f32,
gvz as f32 - (global_size / 2) as f32,
)
.length();
let child_diagonal_length = (child_size as f32 / 2.) * f32::sqrt(3.);
let alpha;
if (dist - 128.).abs() <= child_diagonal_length
if (dist - (global_size as f32 / 2.)).abs() <= child_diagonal_length
{
children.push(StructurePointer::new(depth <= 3, false, 0));
//children.push(StructurePointer::new(depth <= 3, false, 0));
children.push(StructurePointer(
if depth < (self.chunk_power - 1)
{
0xFFFFFFFF
}
else
{
0
},
));
alpha = if dist > 128. { 0. } else { 1. };
}
else if dist > 128.
else if dist > (global_size as f32 / 2.)
{
children.push(StructurePointer::new(false, false, 0));
children.push(StructurePointer(0));
alpha = 0.;
}
else
{
children.push(StructurePointer::new(false, false, 0));
children.push(StructurePointer(0));
alpha = 1.;
}
@@ -86,3 +110,340 @@ where
}
}
}
pub struct SineGenerator<const N: usize>
{
chunk_power: usize,
}
impl<const N: usize> SineGenerator<N>
where
[(); N * N * N]:,
{
pub fn new(chunk_power: usize) -> Self
{
Self { chunk_power }
}
}
impl<const N: usize> Producer<N> for SineGenerator<N>
where
[(); N * N * N]:,
{
fn produce_node(&self, depth: usize, nx: usize, ny: usize, nz: usize) -> ExplicitNTreeNode<N>
{
let node_size = N.pow((self.chunk_power - depth) as u32);
let child_size = node_size / N;
let global_size = N.pow(self.chunk_power as u32);
let mut children = vec![];
let mut children_color = vec![];
let gnx = nx * node_size;
let gny = ny * node_size;
let gnz = nz * node_size;
// Iterate on children of this node
for ((cx, cy), cz) in (0..N).cartesian_product(0..N).cartesian_product(0..N)
{
let total_sub_voxels = child_size * child_size * child_size;
let mut filled_sub_voxels = 0;
let gcx = gnx + cx * child_size;
let gcy = gny + cy * child_size;
let gcz = gnz + cz * child_size;
// prepare 2d values
// Iterate on children
for (x, z) in (0..child_size).cartesian_product(0..child_size)
{
let gvx = gcx + x;
let gvz = gcz + z;
//let gvy = gcy;
let sx = map(gvx as f32, 0., global_size as f32, -8., 8.).abs();
let sz = map(gvz as f32, 0., global_size as f32, -8., 8.).abs();
let sample = (fastapprox::fast::cos(sx) + fastapprox::fast::cos(sz)) * 0.5;
let sample_height = map(sample, -1., 1., 0., 500.);
let prop = map(
sample_height,
gcy as f32,
(gcy + child_size) as f32,
0.,
child_size as f32,
)
.clamp(0., child_size as f32)
.floor() as usize;
filled_sub_voxels += prop;
}
let alpha;
if filled_sub_voxels == 0
{
children.push(StructurePointer(0));
alpha = 0.;
}
else if filled_sub_voxels >= total_sub_voxels
{
children.push(StructurePointer(0));
alpha = 1.;
}
else
{
children.push(StructurePointer(
if depth < (self.chunk_power - 1)
{
0xFFFFFFFF
}
else
{
0
},
));
//children.push(StructurePointer(0));
alpha = if filled_sub_voxels > total_sub_voxels / 2
{
1.
}
else
{
0.
};
}
children_color.push(Color(
(gcx + child_size / 2) as f32 / global_size as f32,
(gcy + child_size / 2) as f32 / global_size as f32,
(gcz + child_size / 2) as f32 / global_size as f32,
alpha,
));
}
ExplicitNTreeNode {
structure: std::array::from_fn(|i| children[i]),
colors: std::array::from_fn(|i| children_color[i]),
}
}
}
pub struct TerrainGenerator<const N: usize>
where
[(); N * N * N]:,
{
chunk_power: usize,
heightmap_width: usize,
heightmap_height: usize,
terrain_width: usize,
terrain_height: usize,
heightmap_min: f32,
heightmap_max: f32,
heightmap: Vec<f32>,
colormap: Vec<u8>,
pub chunk_width: usize,
pub chunk_height: usize,
pub chunk_alt: usize,
}
impl<const N: usize> TerrainGenerator<N>
where
[(); N * N * N]:,
{
pub fn new<P: AsRef<Path>>(
chunk_power: usize,
height_path: P,
height_factor: f32,
color_path: P,
) -> Self
{
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()
{
tiff::decoder::DecodingResult::F32(vec) => vec,
_ => panic!("Unsupported format"),
};
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 terrain_width = color.width() as usize;
let terrain_height = color.height() as usize;
let heightmap_min = heightmap.iter().copied().reduce(f32::min).unwrap();
let heightmap_max = heightmap.iter().copied().reduce(f32::max).unwrap();
// Decide size in chunks
let height_amplitude = heightmap_max - heightmap_min;
let chunk_size = N.pow(chunk_power as u32);
let chunk_width = terrain_width.div_ceil(chunk_size);
let chunk_height = terrain_height.div_ceil(chunk_size);
let chunk_alt = ((height_amplitude / height_factor) as usize).div_ceil(chunk_size);
Self {
chunk_power,
heightmap_width,
heightmap_height,
heightmap_min,
heightmap_max,
terrain_width,
terrain_height,
heightmap,
colormap,
chunk_width,
chunk_height,
chunk_alt,
}
}
}
impl<const N: usize> ChunkedProducer<N> for TerrainGenerator<N>
where
[(); N * N * N]:,
{
fn produce_node(
&self,
depth: usize,
nx: usize,
ny: usize,
nz: usize,
chunk_pos: (usize, usize, usize),
) -> ExplicitNTreeNode<N>
{
let node_size = N.pow((self.chunk_power - depth) as u32);
let child_size = node_size / N;
let global_size = N.pow(self.chunk_power as u32);
let mut children = vec![StructurePointer(0); N * N * N];
let mut children_color = vec![Color(0., 0., 0., 0.); N * N * N];
let gnx = chunk_pos.0 * global_size + nx * node_size;
let gny = chunk_pos.1 * global_size + ny * node_size;
let gnz = chunk_pos.2 * global_size + nz * node_size;
// Iterate on children of this node
for (cx, cz) in (0..N).cartesian_product(0..N)
{
let gcx = gnx + cx * child_size;
let gcz = gnz + cz * child_size;
// prepare 2d values
// Iterate on children
let mut sample_max = self.heightmap_min;
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)
{
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;
}
color_avg.0 /= count as f32;
color_avg.1 /= count as f32;
color_avg.2 /= count as f32;
let sample_min = map(
sample_min,
self.heightmap_min,
self.heightmap_max,
0.,
(self.chunk_alt * global_size) as f32,
) as usize;
let sample_max = map(
sample_max,
self.heightmap_min,
self.heightmap_max,
0.,
(self.chunk_alt * global_size) as f32,
) as usize;
for cy in 0..N
{
let gcy = gny + cy * child_size;
let index = cz * N * N + cy * N + cx;
let alpha;
if gcy > sample_max
{
children[index] = StructurePointer(0);
alpha = 0.;
}
else if gcy + child_size < sample_min
{
children[index] = StructurePointer(0);
alpha = 1.;
}
else
{
children[index] = StructurePointer(
if depth < (self.chunk_power - 1)
{
0xFFFFFFFF
}
else
{
0
},
);
//children.push(StructurePointer(0));
alpha = if (sample_max + sample_min / 2) > gcy + (child_size / 2)
{
1.
}
else
{
0.
};
}
children_color[index] = Color(color_avg.0, color_avg.1, color_avg.2, alpha);
}
}
ExplicitNTreeNode {
structure: std::array::from_fn(|i| children[i]),
colors: std::array::from_fn(|i| children_color[i]),
}
}
}
pub trait ChunkedProducer<const N: usize>
where
[(); N * N * N]:,
{
fn produce_node(
&self,
depth: usize,
nx: usize,
ny: usize,
nz: usize,
chunk_pos: (usize, usize, usize),
) -> ExplicitNTreeNode<N>;
}
src/voxel/cache.rs
+67 -33
View File
@@ -209,32 +209,39 @@ where
@compute
@workgroup_size(16)
@workgroup_size(64)
fn main(
@builtin(global_invocation_id) global_invocation_id: vec3<u32>
)
{{
let request_count_round_up =
((request_count / {children_count}) + select(u32(0), u32(1), request_count % {children_count} != 0)) * {children_count};
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_total_width = u32(1 << (phase + 1));
/*
let phase_count = u32(ceil(log2(f32(request_count))));
if(phase > phase_count)
{{
//return;
return;
}}
*/
/*
if(phase == 0)
{{
let a = index * 2;
@@ -242,6 +249,7 @@ where
sort_indirection[a].child = a % {children_count};
sort_indirection[b].child = b % {children_count};
}}
*/
var swap_indices = vec2<u32>(0, 0);
@@ -255,7 +263,9 @@ where
}}
// Do swap
if(swap_indices.y >= request_count_round_up)
//if(swap_indices.y >= request_count_round_up)
if(swap_indices.y >= total)
{{
// Suppose that swap_indices.y is -inf, dont swap
return;
@@ -306,7 +316,7 @@ where
@group(0) @binding(2) var<storage, read_write> request_counts: u32;
@compute
@workgroup_size(16)
@workgroup_size(64)
fn main(
@builtin(global_invocation_id) global_invocation_id: vec3<u32>
)
@@ -398,7 +408,7 @@ where
@group(0) @binding(2) var<storage, read_write> request_counts: atomic<u32>;
@compute
@workgroup_size(16)
@workgroup_size(64)
fn main(
@builtin(global_invocation_id) global_invocation_id: vec3<u32>
)
@@ -455,7 +465,7 @@ where
@group(0) @binding(2) var<storage, read_write> request_count: atomic<u32>;
@compute
@workgroup_size(16)
@workgroup_size(64)
fn main(
@builtin(global_invocation_id) global_invocation_id: vec3<u32>
)
@@ -469,6 +479,10 @@ 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;
}}
}}
@@ -522,7 +536,7 @@ where
compute_pass.set_pipeline(&self.reset_pipeline);
let shader_invocations = self.cache_size; // one invocation per element
let workgroup_invocations = shader_invocations.div_ceil(16);
let workgroup_invocations = shader_invocations.div_ceil(64);
compute_pass.dispatch_workgroups(workgroup_invocations as u32, 1, 1);
}
@@ -535,8 +549,9 @@ where
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(16);
let workgroups_invocations = request_element_count.div_ceil(64);
/*
// = Perform list compaction
// == Running sum
@@ -562,21 +577,22 @@ where
// == Stream compaction
compute_pass.set_pipeline(&self.compaction_pipeline);
compute_pass.dispatch_workgroups(workgroups_invocations as u32, 1, 1);
*/
// == Bitonic sort
compute_pass.set_pipeline(&self.sort_pipeline);
let sort_element_count = request_element_count * N * N * N;
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
for i in 0..=phases_upper_bound
{
for j in 0..=i
{
compute_pass.set_immediates(0, bytemuck::bytes_of(
&(i | (j << 16))
));
compute_pass.dispatch_workgroups(sort_element_count.div_ceil(2).div_ceil(16) as u32, 1, 1);
compute_pass.dispatch_workgroups(sort_element_count.div_ceil(2).div_ceil(64) as u32, 1, 1);
}
}
}
@@ -714,7 +730,7 @@ impl UsageBuffer
@compute
@workgroup_size(16)
@workgroup_size(64)
fn main(
@builtin(global_invocation_id) global_invocation_id: vec3<u32>
)
@@ -810,9 +826,9 @@ impl UsageBuffer
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);
let workgroup_invocations = shader_invocations.div_ceil(64);
for i in 0..sort_steps
for i in 0..=sort_steps
{
for j in 0..=i
{
@@ -842,8 +858,8 @@ impl StructureTable
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,
mapped_at_creation: false,
usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC,
mapped_at_creation: false,
});
StructureTable {
@@ -867,7 +883,7 @@ impl StructureTable
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,
usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC | BufferUsages::COPY_DST,
mapped_at_creation: false,
});
@@ -986,11 +1002,25 @@ pub struct DestinationElement
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: [Color; N * N * N],
colors: [ColorBytes; N * N * N],
}
pub struct LocationPoolElement
@@ -1343,7 +1373,7 @@ where
struct ColorPoolElement
{{
colors: array<vec4<f32>, {children_count}>
colors: array<u32, {children_count}>
}}
struct LocationPoolElement
@@ -1427,7 +1457,7 @@ where
@group(2) @binding(0) var<storage, read_write> requests: array<CacheNodeRequest>;
@compute
@workgroup_size(16)
@workgroup_size(64)
fn main(
@builtin(global_invocation_id) global_invocation_id: vec3<u32>
)
@@ -1435,7 +1465,7 @@ where
// One shader invocation per invocation on both structure table domain
// and cache domain
var index = global_invocation_id.x;
let max_requests_count = arrayLength(&requests);
let max_requests_count = min(arrayLength(&requests), arrayLength(&lru_list));
total_request_count = min(max_requests_count, pools_request_count + structure_table_request_count);
if(index >= total_request_count)
{{
@@ -1586,7 +1616,7 @@ where
@group(2) @binding(3) var<storage, read> destinations: array<DestinationElement>;
@compute
@workgroup_size(16)
@workgroup_size(64)
fn main(
@builtin(global_invocation_id) global_invocation_id: vec3<u32>
)
@@ -1594,7 +1624,8 @@ where
// Copy with indirection
let index = global_invocation_id.x;
let total = arrayLength(&structure_nodes);
if(index >= total)
let total_cache = arrayLength(&lru_list);
if(index >= total || index >= total_cache)
{{
return;
}}
@@ -1604,7 +1635,10 @@ where
{{
// Phase 1
// Copy into cache page
structure_pool[overwritten_element] = structure_nodes[index];
for(var i = 0; i < {children_count}; i++)
{{
structure_pool[overwritten_element].pointers[i] = select(u32(0), u32(1)<<31, structure_nodes[index].pointers[i] != 0);
}}
color_pool[overwritten_element] = color_nodes[index];
location_pool[overwritten_element] = locations[index];
@@ -1620,7 +1654,7 @@ where
if(destinations[index].child == 0xFFFFFFFF)
{{
structure_table_pointers[destinations[index].node] = new_pointer;
}}else
}}else if usage_buffer[destinations[index].node] != parameters.frame_timestamp + 1
{{
structure_pool[destinations[index].node].pointers[destinations[index].child] = new_pointer;
}}
@@ -1662,7 +1696,7 @@ where
var<immediate> frame_timestamp: u32;
@compute
@workgroup_size(16)
@workgroup_size(64)
fn main(
@builtin(global_invocation_id) global_invocation_id: vec3<u32>
)
@@ -1863,7 +1897,7 @@ where
// Compute necessary shader invocations
let shader_invocation_count = request_target.size() / size_of::<CacheNodeRequest>() as u64;
let workgroup_invocation_count = shader_invocation_count.div_ceil(16);
let workgroup_invocation_count = shader_invocation_count.div_ceil(64);
write_requests_pass.dispatch_workgroups(workgroup_invocation_count as u32, 1, 1);
}
@@ -1930,7 +1964,7 @@ where
// Compute dispatch amounts
let shader_invocation_count =
insertion.structure_nodes.size() as usize / size_of::<StructurePoolElement<N>>();
let workgroup_invocations = shader_invocation_count.div_ceil(16);
let workgroup_invocations = shader_invocation_count.div_ceil(64);
cache_insertion_pass.dispatch_workgroups(workgroup_invocations as u32, 1, 1);
}
@@ -1944,7 +1978,7 @@ where
// Compute dispatch amounts
let shader_invocation_count = self.size + self.structure_table.allocation_table.len();
let workgroup_invocations = shader_invocation_count.div_ceil(16);
let workgroup_invocations = shader_invocation_count.div_ceil(64);
cache_insertion_pass.dispatch_workgroups(workgroup_invocations as u32, 1, 1);
}
@@ -1964,7 +1998,7 @@ where
// Compute dispatch amounts
let shader_invocation_count =
insertion.structure_nodes.size() as usize / size_of::<StructurePoolElement<N>>();
let workgroup_invocations = shader_invocation_count.div_ceil(16);
let workgroup_invocations = shader_invocation_count.div_ceil(64);
cache_insertion_pass.dispatch_workgroups(workgroup_invocations as u32, 1, 1);
}