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

Compare commits

base: octagonal/vxls:z-prepass-experiment
Branches Tags
octagonal/vxls:main octagonal/vxls:z-prepass-experiment
..
compare: octagonal/vxls:a342643ab741eb77fed030badc471a33036b99bd
Branches Tags
octagonal/vxls:main octagonal/vxls:z-prepass-experiment
2 Commits
z-prepass-experiment .. a342643ab7
Author SHA1 Message Date
octagonal a342643ab7 Test lfs 2026-09-11 17:17:22 +02:00
octagonal 54c5e91a7f Fix RT timer 2026-09-11 17:12:38 +02:00
7 changed files with 57 additions and 733 deletions
Expand all files Collapse all files
.gitattributes
+3
View File
@@ -0,0 +1,3 @@
img.jpg filter=lfs diff=lfs merge=lfs -text
img_low.jpg filter=lfs diff=lfs merge=lfs -text
vxls_height.tif filter=lfs diff=lfs merge=lfs -text
Cargo.toml
-1
View File
@@ -16,7 +16,6 @@ glam = "0.33.5"
image = "0.25.10"
indicatif = "0.18.6"
itertools = "0.15.0"
ordered-float = "5.5.0"
pollster = "1.0.1"
rand = "0.10.2"
rayon = "1.12.0"
shaders/voxel.wgsl
+23 -310
View File
@@ -2,12 +2,11 @@ struct VertexOutput
{
@builtin(position) postion: vec4<f32>,
@location(0) @interpolate(flat) chunk_index: u32,
@location(1) ndc: vec4<f32>,
@location(2) color: vec4<f32>,
@location(3) cam_pos: vec3<f32>,
@location(4) world_pos: vec3<f32>,
@location(5) @interpolate(flat) structure_id: u32,
@location(6) chunk_position: vec3<f32>,
@location(1) color: vec4<f32>,
@location(2) cam_pos: vec3<f32>,
@location(3) world_pos: vec3<f32>,
@location(4) @interpolate(flat) structure_id: u32,
@location(5) chunk_position: vec3<f32>
}
struct ChunkImmediate
@@ -15,8 +14,6 @@ struct ChunkImmediate
view_proj: mat4x4<f32>,
cam_pos: vec3<f32>,
frame_timestamp: u32,
width: u32,
downsampling_factor: u32,
}
var<immediate> constants: ChunkImmediate;
@@ -27,7 +24,7 @@ struct CacheChunkObject
transform: mat4x4<f32>,
color: vec4<f32>,
id: u32,
pointer: u32,
pointer: u32
}
@@ -76,7 +73,10 @@ fn unpack_color(color: u32) -> vec4<f32>
@group(0) @binding(5) var<storage, read_write> structure_table_pointer: array<u32>;
@group(0) @binding(6) var<storage, read_write> structure_table_request_buffer: array<atomic<u32>>;
@group(1) @binding(0) var prepass_depth: texture_2d<f32>;
struct FragmentOutput {
@location(0) color: vec4<f32>,
@builtin(frag_depth) depth: f32, // Equivalent to gl_FragDepth
}
@vertex
fn chunk(@builtin(vertex_index) index: u32, @location(0) position: vec3<f32>, @location(1) id: u32) -> VertexOutput
@@ -119,7 +119,6 @@ fn chunk(@builtin(vertex_index) index: u32, @location(0) position: vec3<f32>, @l
var output: VertexOutput;
output.postion = output_vertex;
output.ndc = output_vertex / output_vertex.w;
output.color = vec4(1.);
output.chunk_index = 0;
output.cam_pos = constants.cam_pos;
@@ -217,11 +216,12 @@ fn new_traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_id: u32, dist_of
{
let max_depth = 5;
let dist_offset_voxel = dist_offset * f32(1 << u32(max_depth * 2));
let fovy_deg = 100. / f32(constants.width);
let fovy_deg = 100. / 1920.;
let fovy_rad = (fovy_deg * 3.14) / 180.;
let cone_factor = 1.414 * f32(constants.downsampling_factor) * (tan(fovy_rad / 2.) * 2.); // How many pixels per distance a voxel takes
let cone_factor = tan(fovy_rad / 2.) * 2.;
let st_pointer = structure_table_pointer[root_id];
if (!node_subdivided(st_pointer))
{
@@ -245,11 +245,11 @@ fn new_traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_id: u32, dist_of
}
//var current_node = node_pointer(st_pointer);
// Record usage
var dfs_stack = array<u32, 6>(node_pointer(st_pointer), 0, 0, 0, 0, 0);
var current_depth = 0;
var current_node = dfs_stack[current_depth];
usage_buffer[dfs_stack[current_depth]] = constants.frame_timestamp;
usage_buffer[current_node] = constants.frame_timestamp;
// Start location
//let voxel_dir = select(vec3(-1), vec3(1), ray_dir >= vec3(0.));
@@ -264,7 +264,6 @@ fn new_traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_id: u32, dist_of
var inv_ray_dir = 1. / ray_dir;
var ray_positive = ray_dir > vec3(0.);
var step_dir = select(vec3(-1), vec3(1), ray_positive);
var min_child_size = cone_factor * dist_offset_voxel;
for(var iter = 0; iter < 400; iter ++)
{
@@ -274,18 +273,17 @@ fn new_traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_id: u32, dist_of
var child_pos = (voxel >> vec3(u32(node_shift - 2))) & vec3(3);
var child_index = child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4;
var current_node = dfs_stack[current_depth];
var pointer = structure_pool[current_node].pointers[child_index];
min_child_size = (length(vec3<f32>(voxel) - pos_origin) + dist_offset_voxel) * cone_factor;
let min_child_size = (length(vec3<f32>(voxel) - pos_origin) + dist_offset_voxel) * cone_factor;
while(node_subdivided(pointer) &&
f32(child_size) / 4 > min_child_size
f32(child_size / 4) >= min_child_size
)
{
if(!node_pointer_valid(pointer) && node_subdivided(pointer))
{
// Record request
atomicAdd(&request_buffer[current_node].requests[child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4], 1);
atomicAdd(&request_buffer[dfs_stack[current_depth]].requests[child_index], 1);
break;
}
@@ -295,9 +293,9 @@ fn new_traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_id: u32, dist_of
node_shift = (max_depth - current_depth) * 2;
child_size = 1 << u32(node_shift - 2);
child_pos = (voxel >> vec3(u32(node_shift - 2))) & vec3(3);
child_index = child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4;
current_node = node_pointer(pointer);
dfs_stack[current_depth] = current_node;
child_index = child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4;
pointer = structure_pool[current_node].pointers[child_index];
@@ -312,9 +310,7 @@ fn new_traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_id: u32, dist_of
{
var result: HitResult;
result.color = unpack_color(color);
//result.color = vec4<f32>(f32(iter) / 400.);
//result.hit_pos = (far_t / f32(1 << u32(max_depth * 2))) * ray_dir + ray_origin;
result.hit_pos = (far_t * ray_dir + pos_origin) / f32(1 << u32(max_depth * 2));
result.hit_pos = (far_t / f32(1 << u32(max_depth * 2))) * ray_dir + ray_origin;
return result;
}
@@ -349,7 +345,7 @@ fn new_traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_id: u32, dist_of
let common_depth = ((countLeadingZeros(bit_diffs_lowest) - i32(32 - max_depth * 2)) / 2);
current_depth = common_depth;
//current_node = dfs_stack[current_depth];
current_node = dfs_stack[current_depth];
}
// Iter max color
@@ -359,105 +355,23 @@ fn new_traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_id: u32, dist_of
return result;
}
struct FragmentOutput {
@location(0) color: vec4<f32>,
@builtin(frag_depth) depth: f32, // Equivalent to gl_FragDepth
}
struct FragmentPrepassOutput {
@location(0) depth_prepass: f32, // Equivalent to gl_FragDepth
//@builtin(frag_depth) depth: f32, // Equivalent to gl_FragDepth
}
//fn fragment_prepass(in: VertexOutput) -> @location(0) vec4<f32>
//@early_depth_test(less_equal)
@fragment
fn fragment_prepass(in: VertexOutput) -> FragmentPrepassOutput
{
let ray_dir = normalize(in.world_pos - in.cam_pos);
let interp = box_inter(in.cam_pos, ray_dir, in.chunk_position + vec3(0.), in.chunk_position + vec3(1));
let ray_origin = in.cam_pos + ray_dir * max(interp.x, 0.) - in.chunk_position;
let result = new_traverse(ray_dir, ray_origin, in.structure_id, length(in.cam_pos - (ray_origin + in.chunk_position)));
let clip_pos = constants.view_proj * vec4(result.hit_pos + in.chunk_position, 1.);
let depth = clip_pos.z / clip_pos.w;
var frag_out: FragmentPrepassOutput;
//frag_out.color = result.color;
//frag_out.depth_prepass = in.postion.z;
frag_out.depth_prepass = length(ray_origin + in.chunk_position - in.cam_pos);
//frag_out.depth_prepass = interp.x;
//frag_out.depth_prepass = 100.;
//frag_out.depth = depth;
//frag_out.depth = depth;
//frag_out.depth_prepass = depth;
return frag_out;
}
@early_depth_test(less_equal)
@fragment
fn fragment(in: VertexOutput) -> FragmentOutput
{
let surface_depth = in.postion.z;
let prepass_depth_sample = textureLoad(prepass_depth, vec2<i32>(in.postion.xy), 0).x;
let lin_depth = (100. * 0.01) / (100. - surface_depth * (100. - 0.01));
if lin_depth < prepass_depth_sample || prepass_depth_sample == -1.
{
discard;
}
//frag_out.color = vec4<f32>(2 * 0.01 / (100. + 0.01 - depth * (100. - 0.01)));
let ray_dir = normalize(in.world_pos - in.cam_pos);
let prepass_origin = in.cam_pos + ray_dir * max(prepass_depth_sample - 0.01, 0.);
let interp = box_inter(prepass_origin, ray_dir, in.chunk_position + vec3(0.), in.chunk_position + vec3(1));
let ray_origin = prepass_origin + ray_dir * max(interp.x, 0.) - in.chunk_position;
//let ray_origin = in.cam_pos + ray_dir * (max(0., lin_depth - 0.1)) - in.chunk_position;
//let ray_origin = in.cam_pos + ray_dir * max(interp.x, 0.) - in.chunk_position;
//let ray_origin = in.cam_pos + ray_dir * max(prepass_depth, 0.) - in.chunk_position;
//let space_ro = in.cam_pos + ray_dir * lin_depth;
//let ray_origin = space_ro - in.chunk_position;
let result = new_traverse(ray_dir, ray_origin, in.structure_id, length(in.cam_pos - (ray_origin + in.chunk_position)));
let clip_pos = constants.view_proj * vec4(result.hit_pos + in.chunk_position, 1.);
let depth = clip_pos.z / clip_pos.w;
var frag_out: FragmentOutput;
frag_out.color = result.color;
//frag_out.color = vec4<f32>(vec3<f32>(prepass_depth_sample) / 100., 1.);
//frag_out.color = vec4<f32>(smpl);
//frag_out.color = vec4<f32>(ray_origin, 1.);
//frag_out.color = result.color;
//frag_out.color = vec4<f32>(ray_origin, 1.);
//frag_out.color = vec4<f32>(vec3<f32>(lin_depth) / 100., 1.);
frag_out.depth = depth;
return frag_out;
//return vec4<f32>(ray_origin, 1.);
//return frag_out;
//return vec4(interp.y / 10.);
}
@fragment
fn _fragment(in: VertexOutput) -> FragmentOutput
{
//frag_out.color = vec4<f32>(2 * 0.01 / (100. + 0.01 - depth * (100. - 0.01)));
let ray_dir = normalize(in.world_pos - in.cam_pos);
let interp = box_inter(in.cam_pos - in.chunk_position, ray_dir, vec3(0.), vec3(1));
let ray_origin = (in.cam_pos - in.chunk_position) + ray_dir * (max(0., interp.x));
let result = new_traverse(ray_dir, ray_origin, in.structure_id, length(in.cam_pos - (ray_origin + in.chunk_position)));
let clip_pos = constants.view_proj * vec4(result.hit_pos + in.chunk_position, 1.);
let depth = clip_pos.z / clip_pos.w;
var frag_out: FragmentOutput;
//frag_out.color = result.color;
frag_out.color = result.color;
//frag_out.color = vec4<f32>(smpl);
//frag_out.color = vec4<f32>((2. * 0.01 * 100.) / (0.01 + 100. - prepass_depth * (100. - 0.01)));
//frag_out.depth = depth;
frag_out.depth = depth;
return frag_out;
//return vec4<f32>(ray_origin, 1.);
@@ -482,204 +396,3 @@ fn fragment(in: VertexOutput) -> @location(0) vec4<f32>
}
*/
/*
fn traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_id: u32, dist_offset: f32) -> vec4<f32>
{
let st_pointer = structure_table_pointer[root_id];
if (!node_subdivided(st_pointer))
{
discard;
return vec4(0., 1., 0., 1.);
}
if(!node_pointer_valid(st_pointer))
{
atomicAdd(&structure_table_request_buffer[root_id], 1);
discard;
return vec4(0., 1., 0., 1.);
}
let fovy_deg = 100.;
let fovy = 3.14159 * (fovy_deg / 180.);
let definition = 1920.;
let cone_fovy = fovy / definition;
let factor = 1.;
let cone_size_factor = 2. * tan(cone_fovy) * factor;
// Current depth of the node we are exploring
var current_depth = 0;
// Index of the current node's data
var current_node = u32(st_pointer & 0x3FFFFFFF);
usage_buffer[current_node] = constants.frame_timestamp;
var dfs_stack = array<u32, 6>(current_node, 0, 0, 0, 0, 0);
// Lut of the node_size per depth
var node_size_lut = array<i32, 6>(
4 * 4 * 4 * 4 * 4,
4 * 4 * 4 * 4,
4 * 4 * 4,
4 * 4,
4,
1,
);
let local_dist_offset = dist_offset * f32(node_size_lut[0]);
// Current node size
var node_size = node_size_lut[0]; // 128
// Size of a child of this node
var child_size = node_size / 4;
// Simple FVT
let t_off = abs(1. / ray_dir);
// Start location
let voxel_dir = select(vec3(-1), vec3(1), ray_dir >= vec3(0.));
var pos_origin = clamp(ray_origin * f32(node_size), vec3(0.), vec3(f32(node_size) - 1.));
var voxel = vec3<i32>(pos_origin);
var last_voxel = voxel;
let wall_offset = select(vec3(0), vec3(1), ray_dir > vec3(0.));
let max_depth = u32(5);
var adaptive_depth = i32(max_depth);
var far_t = 0.;
let ray_dir_inv = 1. / ray_dir;
let fma_offset = - pos_origin * ray_dir_inv;
//let depth_limit = 3;
for(var iter = 0; iter < 400; iter ++)
{
// Our ray is currently touching a voxel.
// Descend to the lowest node that contains this voxel
// Position of the child we are in
var child_pos = (vec3<u32>(voxel) >> vec3<u32>((max_depth - u32(current_depth + 1)) * 2)) & vec3<u32>(3); // Hardcode for 4-tree
var child_index = child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4;
// Current node has been used. report
usage_buffer[current_node] = constants.frame_timestamp;
while(
node_subdivided(structure_pool[current_node].pointers[child_index]) &&
(local_dist_offset + far_t) * cone_size_factor < f32(node_size_lut[current_depth])
)
{
if(!node_pointer_valid(structure_pool[current_node].pointers[child_index]))
{
atomicAdd(&request_buffer[current_node].requests[child_index], 1);
break;
}
// Child node is subdivided, we go in, save position in stack
current_node = node_pointer(structure_pool[current_node].pointers[child_index]);
usage_buffer[current_node] = constants.frame_timestamp;
current_depth += 1;
dfs_stack[current_depth] = current_node;
node_size = node_size_lut[current_depth];
child_pos = (vec3<u32>(voxel) >> vec3<u32>((max_depth - u32(current_depth + 1)) * 2)) & vec3<u32>(3); // Hardcode for 4-tree
child_index = child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4;
child_size = node_size / 4;
}
usage_buffer[current_node] = constants.frame_timestamp;
// At this point current_depth is the depth of the node that contains the voxel
// child_pos and child_index relate to the specific child of the node that contains this voxel
// It is guaranteed that the child is leave
// Check current leave's color
let color = unpack_color(color_pool[current_node].colors[child_index]);
if(color.w != 0.) // Not transparent
{
/*
let k = child_pos.x + child_pos.y + child_pos.z;
let w = voxel.x + voxel.y + voxel.z;
let x = select(0.5, 1., k % 2 == 0) * select(0.8, 1., w % 2 == 0);
var div = 1;
var overlay = 1.;
for(var i = 1; i <= 5; i++)
{
let x = (voxel.x / div + voxel.y / div + voxel.z / div) % 2 == 0;
overlay -= select(0., 1. / (f32(i) * 2.5), x);
div *= 4;
}
*/
return color;
//return overlay * color;
}
// Voxel and whole child containing it is empty
// Perform a step through the children of the node
let child_position = (voxel / child_size) * child_size;
let far_corner = child_position + wall_offset * child_size;
//let far_ts = (vec3<f32>(far_corner) - pos_origin) / ray_dir; // TODO: Turn into fma
let far_ts = fma(vec3<f32>(far_corner), ray_dir_inv, fma_offset);
far_t = min(min(far_ts.x, far_ts.y), far_ts.z);
let next_child_min = select(child_position, child_position + voxel_dir * child_size, vec3(far_t) == far_ts);
let next_child_max = next_child_min + vec3(child_size) - vec3(1);
// The ray (far_t) is now touching the new child to explore
// Find out which actual voxel we are touching
let previous_voxel = voxel;
let float_voxel = clamp(vec3<i32>(pos_origin + far_t * ray_dir), next_child_min, next_child_max);
/*
voxel = vec3<i32>(
floor(
select(
float_voxel - vec3(0.5),
float_voxel + vec3(0.5),
ray_dir > vec3(0.)
))
);
*/
//voxel = vec3<i32>(round(float_voxel));
//voxel = voxel_from_wall(float_voxel, ray_dir);
//voxel = voxel_from_wall(float_voxel, ray_dir);
voxel = float_voxel;
if(any(voxel < vec3(0)) || any(voxel >= vec3(node_size_lut[0])))
{
//return vec4(f32(iter) / 100.);
discard;
}
// We touched a voxel as if we explored blocks sized by the child size of the current node.
// But we might have exited the current node.
// If this is the case we have to walk back up the tree
// And then back down to the next node over
// As such we find the lowest ancestor that can contain both the privous voxel (in node) and the new voxel (out of node)
let bit_diffs = voxel ^ previous_voxel;
let bit_diffs_lowest = bit_diffs.x | bit_diffs.y | bit_diffs.z;
let flb = ((countLeadingZeros(bit_diffs_lowest) - i32(32 - max_depth * 2)) / 2);
let common_depth = flb;
current_depth = common_depth;
node_size = node_size_lut[current_depth];
child_size = node_size / 4;
current_node = dfs_stack[current_depth];
// Figure out current voxel position
//voxel = vec3<i32>(ray_origin + ray_dir * t);
}
return vec4<f32>(1., 0., 1., 1.);
}
*/
src/host_production.rs
-15
View File
@@ -1,15 +0,0 @@
// Contains useful facilities to render data streamed in from the host
// Can produce a voxel given
// - Its depth
// - Its position within the chunk
// - The chunks position
pub trait ChunkVoxelProducer
{
fn produce_voxel(
&mut self,
depth: usize,
chunk_position: (usize, usize, usize),
voxel_position: (usize, usize, usize),
);
}
src/main.rs
+27 -402
View File
@@ -16,16 +16,11 @@ use glam::Mat4;
use glam::Vec3;
use glam::Vec4;
use itertools::Itertools;
use ordered_float::OrderedFloat;
use rayon::iter::IndexedParallelIterator;
use rayon::iter::IntoParallelRefIterator;
use rayon::iter::ParallelIterator;
use wgpu::BindGroupLayout;
use wgpu::BlendState;
use wgpu::Buffer;
use wgpu::BufferUsages;
use wgpu::Color;
use wgpu::ComputePipeline;
use wgpu::Device;
use wgpu::Extent3d;
use wgpu::Features;
@@ -33,12 +28,8 @@ use wgpu::InstanceDescriptor;
use wgpu::InstanceFlags;
use wgpu::MemoryBudgetThresholds;
use wgpu::Operations;
use wgpu::Origin3d;
use wgpu::PipelineCompilationOptions;
use wgpu::RenderPipeline;
use wgpu::ShaderStages;
use wgpu::Texture;
use wgpu::TextureFormat;
use wgpu::TextureUsages;
use wgpu::TextureView;
use wgpu::util::BufferInitDescriptor;
@@ -53,7 +44,6 @@ 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;
@@ -74,7 +64,6 @@ use crate::voxel_cache::producer_interface::CacheRequest;
mod camera;
mod egui_renderer;
mod host_production;
mod producers;
mod sparse_tree;
mod voxel_cache;
@@ -89,18 +78,10 @@ struct State
size: winit::dpi::PhysicalSize<u32>,
surface: wgpu::Surface<'static>,
depth_buffer: (wgpu::Texture, wgpu::TextureView),
prepass_depth_buffer: (wgpu::Texture, wgpu::TextureView),
prepass_depth: (wgpu::Texture, wgpu::TextureView),
upsampled_prepass_depth: (wgpu::Texture, wgpu::TextureView),
prepass_downsampling: u32,
upsample_pipeline: ComputePipeline,
surface_format: wgpu::TextureFormat,
egui_renderer: EguiRenderer,
pipeline: RenderPipeline,
prepass_pipeline: RenderPipeline,
prepass_upsample_bg_layout: BindGroupLayout,
prepass_depth_bind_group_layout: BindGroupLayout,
voxel_cache: Arc<Mutex<VoxelCache<4>>>,
cache_interface: Arc<CacheProducerInterface<4>>,
terrain_generator: Arc<TerrainGenerator<4>>,
@@ -126,8 +107,6 @@ struct Immediates
view_proj: Mat4,
cam_pos: Vec3,
frame_timestamp: u32,
width: u32,
downsample_factor: u32,
}
#[derive(Debug, Clone, Copy, Zeroable, Pod)]
@@ -179,11 +158,10 @@ impl State
let egui_renderer = EguiRenderer::new(&device, surface_format, &window);
let mut voxel_cache = VoxelCache::<4>::new(100_000, device.clone(), queue.clone());
let mut voxel_cache = VoxelCache::<4>::new(200_000, device.clone(), queue.clone());
let cache_interface = CacheProducerInterface::new(1024, &device);
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",
@@ -192,9 +170,9 @@ impl State
// );
// let terrain_generator = TerrainGenerator::<4>::new(
// 5,
// "/home/albin/Documents/vxls_maps/orgere/height.tif",
// "/home/albin/Documents/vxls_maps/lapiz/height.tif",
// 0.2,
// "/home/albin/Documents/vxls_maps/orgere/color.jpg",
// "/home/albin/Documents/vxls_maps/lapiz/ortho.jpg",
// );
let mut chunk_pos_map = HashMap::new();
@@ -231,98 +209,16 @@ impl State
),
});
let prepass_pipeline_layout =
device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
label: Some("Voxel pipeline layout"),
let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
label: Some("Voxel pipeline layout"),
bind_group_layouts: &[Some(&voxel_cache.bind_group_layout())],
immediate_size: size_of::<Immediates>() as u32,
});
let prepass_pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
label: Some("Render pipeline"),
layout: Some(&prepass_pipeline_layout),
vertex: wgpu::VertexState {
module: &shader_module,
entry_point: Some("chunk"),
compilation_options: Default::default(),
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: Some(wgpu::Face::Front),
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::Less),
stencil: wgpu::StencilState::default(),
bias: wgpu::DepthBiasState::default(),
}),
multisample: wgpu::MultisampleState::default(),
fragment: Some(wgpu::FragmentState {
module: &shader_module,
entry_point: Some("fragment_prepass"),
compilation_options: wgpu::PipelineCompilationOptions::default(),
targets: &[Some(wgpu::ColorTargetState {
format: TextureFormat::R32Float,
blend: None,
write_mask: wgpu::ColorWrites::all(),
})],
}),
multiview_mask: None,
cache: None,
bind_group_layouts: &[Some(&voxel_cache.bind_group_layout())],
immediate_size: size_of::<Immediates>() as u32,
});
let prepass_depth_bind_group_layout =
device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
label: Some("prepass_upsample_bind_group_layout "),
entries: &[wgpu::BindGroupLayoutEntry {
binding: 0,
visibility: ShaderStages::FRAGMENT,
ty: wgpu::BindingType::Texture {
sample_type: wgpu::TextureSampleType::Float { filterable: false },
view_dimension: wgpu::TextureViewDimension::D2,
multisampled: false,
},
count: None,
}],
});
let chunk_pipeline_layout =
device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
label: Some("Voxel pipeline layout"),
bind_group_layouts: &[
Some(&voxel_cache.bind_group_layout()),
Some(&prepass_depth_bind_group_layout),
],
immediate_size: size_of::<Immediates>() as u32,
});
let chunk_pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
label: Some("Render pipeline"),
layout: Some(&chunk_pipeline_layout),
layout: Some(&pipeline_layout),
vertex: wgpu::VertexState {
module: &shader_module,
entry_point: Some("chunk"),
@@ -375,85 +271,6 @@ impl State
cache: None,
});
let prepass_upsample_bg_layout =
device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
label: Some("prepass_upsample_bg"),
entries: &[
wgpu::BindGroupLayoutEntry {
binding: 0,
visibility: ShaderStages::COMPUTE,
ty: wgpu::BindingType::StorageTexture {
access: wgpu::StorageTextureAccess::ReadOnly,
format: wgpu::TextureFormat::R32Float,
view_dimension: wgpu::TextureViewDimension::D2,
},
count: None,
},
wgpu::BindGroupLayoutEntry {
binding: 1,
visibility: ShaderStages::COMPUTE,
ty: wgpu::BindingType::StorageTexture {
access: wgpu::StorageTextureAccess::WriteOnly,
format: wgpu::TextureFormat::R32Float,
view_dimension: wgpu::TextureViewDimension::D2,
},
count: None,
},
],
});
let prepass_upsample_layout =
device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
label: Some("prepass_upsample_layout"),
bind_group_layouts: &[Some(&prepass_upsample_bg_layout)],
immediate_size: 0,
});
let prepass_downsampling = 4;
let prepass_upsample = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
label: Some("prepass_upsample"),
layout: Some(&prepass_upsample_layout),
module: &device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: Some("prepass_upsample_module"),
source: wgpu::ShaderSource::Wgsl(
format!(
"
@group(0) @binding(0) var input_tex: texture_storage_2d<r32float, read>;
@group(0) @binding(1) var output_tex: texture_storage_2d<r32float, write>;
@compute @workgroup_size(8, 8)
fn main(@builtin(global_invocation_id) id: vec3<u32>)
{{
let source_pos = vec2<i32>(id.xy);
var depth = textureLoad(input_tex, source_pos + vec2<i32>(0, 0)).x;
depth = min(depth, textureLoad(input_tex, source_pos + vec2<i32>(1, 0)).x);
depth = min(depth, textureLoad(input_tex, source_pos + vec2<i32>(0, 1)).x);
depth = min(depth, textureLoad(input_tex, source_pos + vec2<i32>(1, 1)).x);
let dest_pos = source_pos * {prepass_downsampling};
for(var ox = 0; ox < {prepass_downsampling}; ox ++)
{{
for(var oy = 0; oy < {prepass_downsampling}; oy ++)
{{
textureStore(
output_tex,
dest_pos + vec2<i32>(ox, oy),
vec4<f32>(depth)
);
}}
}}
}}
"
)
.into(),
),
}),
entry_point: Some("main"),
compilation_options: PipelineCompilationOptions::default(),
cache: None,
});
let state = State {
instance,
window,
@@ -462,33 +279,13 @@ impl State
surface_format,
egui_renderer,
depth_buffer: Self::create_depth_buffer(&device, size.width, size.height),
prepass_depth_buffer: Self::create_depth_buffer(
&device,
size.width / prepass_downsampling,
size.height / prepass_downsampling,
),
prepass_depth: Self::create_prepass_depth_buffer(
&device,
size.width / prepass_downsampling,
size.height / prepass_downsampling,
),
upsampled_prepass_depth: Self::create_prepass_depth_buffer(
&device,
size.width,
size.height,
),
prepass_downsampling,
upsample_pipeline: prepass_upsample,
queue,
device,
usage_vec: Arc::new(Mutex::new(vec![])),
pipeline: chunk_pipeline,
prepass_pipeline,
prepass_upsample_bg_layout,
prepass_depth_bind_group_layout,
voxel_cache: Arc::new(Mutex::new(voxel_cache)),
cache_interface: cache_interface.into(),
insertion_debounce: true,
insertion_debounce: false,
camera: Default::default(),
instance_buffer,
instance_count,
@@ -510,19 +307,6 @@ impl State
fn handle_event(&mut self, event: &WindowEvent)
{
if let WindowEvent::KeyboardInput { event, .. } = event
{
match (event.state, event.physical_key)
{
(
winit::event::ElementState::Pressed,
winit::keyboard::PhysicalKey::Code(winit::keyboard::KeyCode::KeyF),
) => self.insertion_debounce = !self.insertion_debounce,
_ =>
{}
}
}
self.egui_renderer.handle_input(&self.window, event);
self.camera.handle_input(event);
}
@@ -542,35 +326,6 @@ impl State
}
}
fn create_prepass_depth_buffer(
device: &Device,
width: u32,
height: u32,
) -> (Texture, TextureView)
{
let texture = device.create_texture(&wgpu::wgt::TextureDescriptor {
label: Some("Prepass Depth buffer"),
size: Extent3d {
width,
height,
depth_or_array_layers: 1,
},
mip_level_count: 1,
sample_count: 1,
dimension: wgpu::TextureDimension::D2,
format: wgpu::TextureFormat::R32Float,
usage: TextureUsages::RENDER_ATTACHMENT
| TextureUsages::TEXTURE_BINDING
| TextureUsages::STORAGE_BINDING,
view_formats: &[wgpu::TextureFormat::R32Float],
});
let texture_view = texture.create_view(&wgpu::wgt::TextureViewDescriptor {
label: Some("prepass depth view"),
..Default::default()
});
(texture, texture_view)
}
fn create_depth_buffer(device: &Device, width: u32, height: u32) -> (Texture, TextureView)
{
let texture = device.create_texture(&wgpu::wgt::TextureDescriptor {
@@ -620,18 +375,6 @@ impl State
self.configure_surface();
self.depth_buffer =
Self::create_depth_buffer(&self.device, new_size.width, new_size.height);
self.prepass_depth_buffer = Self::create_depth_buffer(
&self.device,
new_size.width / self.prepass_downsampling,
new_size.height / self.prepass_downsampling,
);
self.prepass_depth = Self::create_prepass_depth_buffer(
&self.device,
new_size.width / self.prepass_downsampling,
new_size.height / self.prepass_downsampling,
);
self.upsampled_prepass_depth =
Self::create_prepass_depth_buffer(&self.device, new_size.width, new_size.height);
}
fn render(&mut self)
@@ -639,43 +382,6 @@ impl State
self.camera.update();
self.voxel_cache.lock().next_frame();
// Build sorted buffer
let mut chunks = self.chunk_pos_map.iter().collect::<Vec<_>>();
chunks.sort_by_key(|(_, (x, y, z))| {
OrderedFloat(
(Vec3::new(*x as f32, *y as f32, *z as f32) - self.camera.position).length(),
)
});
let instance_buffer = self
.device
.create_buffer_init(&wgpu::util::BufferInitDescriptor {
label: Some("Instance buffer"),
contents: bytemuck::cast_slice(
chunks
.iter()
.map(|(structure_id, (x, y, z))| InstanceAttribute {
x: *x as f32,
y: *y as f32,
z: *z as f32,
id: **structure_id,
})
.collect::<Vec<_>>()
.as_slice(),
),
usage: BufferUsages::COPY_DST | BufferUsages::VERTEX,
});
let prepass_depth_bind_group = self.device.create_bind_group(&wgpu::BindGroupDescriptor {
label: Some("prepass_depth_bind_group"),
layout: &self.prepass_depth_bind_group_layout,
entries: &[wgpu::BindGroupEntry {
binding: 0,
resource: wgpu::BindingResource::TextureView(&self.upsampled_prepass_depth.1),
}],
});
// Create texture view.
// NOTE: We must handle Timeout because the surface may be unavailable
// (e.g., when the window is occluded on macOS).
@@ -730,88 +436,8 @@ impl State
mapped_at_creation: false,
});
// ~~ Prepass upsample bind group ~~
let prepass_bg = self.device.create_bind_group(&wgpu::BindGroupDescriptor {
label: Some("prepass_bind_group"),
layout: &self.prepass_upsample_bg_layout,
entries: &[
wgpu::BindGroupEntry {
binding: 0,
resource: wgpu::BindingResource::TextureView(&self.prepass_depth.1),
},
wgpu::BindGroupEntry {
binding: 1,
resource: wgpu::BindingResource::TextureView(&self.upsampled_prepass_depth.1),
},
],
});
let mut encoder = self.device.create_command_encoder(&Default::default());
// ~~ Prepass ~~
{
let mut renderpass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
label: None,
color_attachments: &[Some(wgpu::RenderPassColorAttachment {
view: &self.prepass_depth.1,
depth_slice: None,
resolve_target: None,
ops: Operations {
load: wgpu::LoadOp::Clear(Color {
r: -1.,
g: 0.,
b: 0.,
a: 0.,
}),
store: wgpu::StoreOp::Store,
},
})],
depth_stencil_attachment: Some(wgpu::RenderPassDepthStencilAttachment {
view: &self.prepass_depth_buffer.1,
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_vertex_buffer(0, instance_buffer.slice(..));
renderpass.set_bind_group(0, Some(&self.voxel_cache.lock().bind_group()), &[]);
let imm = [Immediates {
view_proj: self.camera.view_proj(),
cam_pos: self.camera.position,
frame_timestamp: self.voxel_cache.lock().current_timestamp(),
width: self.size.width / self.prepass_downsampling,
downsample_factor: self.prepass_downsampling,
}];
renderpass.set_pipeline(&self.prepass_pipeline);
renderpass.set_immediates(0, unsafe { as_raw_bytes(&imm) });
renderpass.draw(0..36, 0..(self.instance_count as u32));
// End the renderpass.
drop(renderpass);
}
// ~~ Upsample pass ~~
{
let mut compute_pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
label: Some("upsample_compute_pass"),
timestamp_writes: None,
});
compute_pass.set_bind_group(0, Some(&prepass_bg), &[]);
compute_pass.set_pipeline(&self.upsample_pipeline);
compute_pass.dispatch_workgroups(
(self.size.width / self.prepass_downsampling).div_ceil(8),
(self.size.height / self.prepass_downsampling).div_ceil(8),
1,
);
}
// ~~ Main render pass ~~
let mut encoder = self.device.create_command_encoder(&Default::default());
{
let mut renderpass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
label: None,
@@ -845,18 +471,15 @@ impl State
occlusion_query_set: None,
multiview_mask: None,
});
renderpass.set_vertex_buffer(0, instance_buffer.slice(..));
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()), &[]);
renderpass.set_bind_group(1, Some(&prepass_depth_bind_group), &[]);
let imm = [Immediates {
view_proj: self.camera.view_proj(),
cam_pos: self.camera.position,
frame_timestamp: self.voxel_cache.lock().current_timestamp(),
width: self.size.width,
downsample_factor: 1,
}];
renderpass.set_pipeline(&self.pipeline);
renderpass.set_immediates(0, unsafe { as_raw_bytes(&imm) });
renderpass.draw(0..36, 0..(self.instance_count as u32));
@@ -869,10 +492,13 @@ impl State
// ~~ EGUI Render pass ~~
{
self.egui_renderer.begin_frame(&self.window);
egui::Window::new("Window ! ").resizable(false).show(
egui::Window::new("Window ! ").resizable(true).show(
self.egui_renderer.context(),
|ui| {
if !self.insertion_debounce
if self
.camera
.pressed_keyset
.contains(&winit::keyboard::KeyCode::KeyF)
{
ui.label(
egui::RichText::new("Cache paused")
@@ -936,11 +562,10 @@ impl State
);
// ~~ Do cache managment
// if !self
// .camera
// .pressed_keyset
// .contains(&winit::keyboard::KeyCode::KeyF)
if self.insertion_debounce
if !self
.camera
.pressed_keyset
.contains(&winit::keyboard::KeyCode::KeyF)
{
self.voxel_cache
.lock()
@@ -968,12 +593,12 @@ impl State
);
// ~~ Do cache managment
// if !self
// .camera
// .pressed_keyset
// .contains(&winit::keyboard::KeyCode::KeyF)
if self.insertion_debounce
if !self
.camera
.pressed_keyset
.contains(&winit::keyboard::KeyCode::KeyF)
{
self.insertion_debounce = true;
let request_count = self
.cache_interface
.total_request_count(&self.device, &self.queue);
src/producers.rs
+1 -3
View File
@@ -3,8 +3,6 @@ use std::path::Path;
use glam::Vec3;
use itertools::Itertools;
use rayon::iter::IntoParallelRefMutIterator;
use rayon::iter::ParallelIterator;
use crate::sparse_tree::Color;
use crate::voxel_cache::data::ExplicitNTreeNode;
@@ -312,7 +310,7 @@ where
let heightmap_max = heightmap.iter().copied().reduce(f32::max).unwrap();
heightmap
.par_iter_mut()
.iter_mut()
.filter(|x| **x == -9999.)
.for_each(|x| *x = heightmap_min);
src/voxel_cache/structure_table.rs
+3 -2
View File
@@ -80,7 +80,7 @@ impl StructureTable
&mut self.request_buffer
}
pub fn free_structure(&mut self, structure_id: u32)
pub fn remove_structure(&mut self, structure_id: u32)
{
self.available_slots += 1;
self.allocation_table[structure_id as usize] = false;
@@ -104,7 +104,8 @@ impl StructureTable
.allocation_table
.iter()
.enumerate()
.find(|(_, allocated)| !**allocated)
.filter(|(_, allocated)| !**allocated)
.next()
.unwrap();
self.allocation_table[first_id] = true;
self.available_slots -= 1;