working z prepass shit

Cargo.toml

@@ -16,6 +16,7 @@ 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

@@ -2,11 +2,12 @@ struct VertexOutput
 {
     @builtin(position) postion: vec4<f32>,
     @location(0) @interpolate(flat) chunk_index: u32,
-    @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>
+    @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>,
 }
 
 struct ChunkImmediate
@@ -14,6 +15,8 @@ struct ChunkImmediate
     view_proj: mat4x4<f32>,
     cam_pos: vec3<f32>,
     frame_timestamp: u32,
+    width: u32,
+    downsampling_factor: u32,
 }
 
 var<immediate> constants: ChunkImmediate;
@@ -24,7 +27,7 @@ struct CacheChunkObject
     transform: mat4x4<f32>,
     color: vec4<f32>,
     id: u32,
-    pointer: u32
+    pointer: u32,
 }
 
 
@@ -73,10 +76,7 @@ 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>>;
 
-struct FragmentOutput {
-    @location(0) color: vec4<f32>,
-    @builtin(frag_depth) depth: f32, // Equivalent to gl_FragDepth
-}
+@group(1) @binding(0) var prepass_depth: texture_2d<f32>;
 
 @vertex
 fn chunk(@builtin(vertex_index) index: u32, @location(0) position: vec3<f32>, @location(1) id: u32) -> VertexOutput
@@ -119,6 +119,7 @@ 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;
@@ -216,15 +217,15 @@ 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. / 1920.;
+    let fovy_deg = 100. / f32(constants.width);
     let fovy_rad = (fovy_deg * 3.14) / 180.;
-    let cone_factor = tan(fovy_rad / 2.) * 2.;
+    let cone_factor = 1.414 * f32(constants.downsampling_factor) * (tan(fovy_rad / 2.) * 2.); // How many pixels per distance a voxel takes
 
     let st_pointer = structure_table_pointer[root_id];
-
     
     if (!node_subdivided(st_pointer))
     {
+        discard;
         var result: HitResult;
         result.color = vec4(0., 1., 0., 1.);
         result.hit_pos = ray_origin;
@@ -236,6 +237,7 @@ fn new_traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_id: u32, dist_of
         // Send request on structure table
         atomicAdd(&structure_table_request_buffer[root_id], 1); 
 
+        discard;
         var result: HitResult;
         result.color = vec4(0., 1., 0., 1.);
         result.hit_pos = ray_origin;
@@ -262,6 +264,7 @@ 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 ++)
     {
@@ -270,17 +273,19 @@ fn new_traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_id: u32, dist_of
         child_size = 1 << u32(node_shift - 2);
 
         var child_pos = (voxel >> vec3(u32(node_shift - 2))) & vec3(3);
-        var pointer = structure_pool[dfs_stack[current_depth]].pointers[child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4];
+        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];
         
-        let min_child_size = (length(vec3<f32>(voxel) - pos_origin) + dist_offset_voxel) * cone_factor;
+        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[dfs_stack[current_depth]].requests[child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4], 1);
+                atomicAdd(&request_buffer[current_node].requests[child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4], 1);
                 break;
             }
 
@@ -290,22 +295,26 @@ 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);
-            dfs_stack[current_depth] = node_pointer(pointer);
+            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;
 
-            pointer = structure_pool[dfs_stack[current_depth]].pointers[child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4];
+            pointer = structure_pool[current_node].pointers[child_index];
 
             // Record usage
-            usage_buffer[dfs_stack[current_depth]] = constants.frame_timestamp;
+            usage_buffer[current_node] = constants.frame_timestamp;
         }
 
 
         // Check color
-        let color = color_pool[dfs_stack[current_depth]].colors[child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4];
+        let color = color_pool[current_node].colors[child_index];
         if(((color >> 24) & 0xFF) != 0)
         {
             var result: HitResult;
             result.color = unpack_color(color);
-            result.hit_pos = (far_t / f32(1 << u32(max_depth * 2))) * ray_dir + ray_origin;
+            //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));
             return result;
         }
 
@@ -350,17 +359,143 @@ 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;
+    return frag_out;
+
+    //return vec4<f32>(ray_origin, 1.);
+    //return frag_out;
+    //return vec4(interp.y / 10.);
+}
+
+/*
+@fragment
+fn fragment(in: VertexOutput) -> @location(0) vec4<f32>
+{
+    let st = structure_table_pointer[0];
+    let subdivided = ((st >> 31) & 1) != 0;
+    let pointer_valid = ((st >> 30) & 1) != 0;
+    // Request stuff
+    atomicAdd(&structure_table_request_buffer[0], 1); 
+    if(subdivided && !pointer_valid)
+    {
+        return vec4(0., 1., 0., 1.);
+    }
+    return vec4(1., 0., 0., 1.);
+}
+*/
+
+/*
 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.);
     }
 
@@ -547,45 +682,4 @@ fn traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_id: u32, dist_offset
     return vec4<f32>(1., 0., 1., 1.);
 
 }
-
-@early_depth_test(less_equal)
-@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.depth = depth;
-    return frag_out;
-
-    //return vec4<f32>(ray_origin, 1.);
-    //return frag_out;
-    //return vec4(interp.y / 10.);
-}
-
-/*
-@fragment
-fn fragment(in: VertexOutput) -> @location(0) vec4<f32>
-{
-    let st = structure_table_pointer[0];
-    let subdivided = ((st >> 31) & 1) != 0;
-    let pointer_valid = ((st >> 30) & 1) != 0;
-    // Request stuff
-    atomicAdd(&structure_table_request_buffer[0], 1); 
-    if(subdivided && !pointer_valid)
-    {
-        return vec4(0., 1., 0., 1.);
-    }
-    return vec4(1., 0., 0., 1.);
-}
 */
-

src/host_production.rs

@@ -0,0 +1,15 @@
+// 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

@@ -16,11 +16,16 @@ 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;
@@ -28,8 +33,12 @@ 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;
@@ -44,6 +53,7 @@ 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;
 
@@ -64,6 +74,7 @@ use crate::voxel_cache::producer_interface::CacheRequest;
 
 mod camera;
 mod egui_renderer;
+mod host_production;
 mod producers;
 mod sparse_tree;
 mod voxel_cache;
@@ -78,10 +89,18 @@ 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>>,
@@ -107,6 +126,8 @@ struct Immediates
     view_proj: Mat4,
     cam_pos: Vec3,
     frame_timestamp: u32,
+    width: u32,
+    downsample_factor: u32,
 }
 
 #[derive(Debug, Clone, Copy, Zeroable, Pod)]
@@ -158,10 +179,11 @@ impl State
 
         let egui_renderer = EguiRenderer::new(&device, surface_format, &window);
 
-        let mut voxel_cache = VoxelCache::<4>::new(200_000, device.clone(), queue.clone());
+        let mut voxel_cache = VoxelCache::<4>::new(100_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",
@@ -170,9 +192,9 @@ impl State
         // );
         // let terrain_generator = TerrainGenerator::<4>::new(
         //     5,
-        //     "/home/albin/Documents/vxls_maps/lapiz/height.tif",
+        //     "/home/albin/Documents/vxls_maps/orgere/height.tif",
         //     0.2,
-        //     "/home/albin/Documents/vxls_maps/lapiz/ortho.jpg",
+        //     "/home/albin/Documents/vxls_maps/orgere/color.jpg",
         // );
 
         let mut chunk_pos_map = HashMap::new();
@@ -209,16 +231,98 @@ impl State
             ),
         });
 
-        let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
-            label: Some("Voxel pipeline layout"),
+        let prepass_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,
+                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,
         });
 
+        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(&pipeline_layout),
+            layout: Some(&chunk_pipeline_layout),
             vertex: wgpu::VertexState {
                 module: &shader_module,
                 entry_point: Some("chunk"),
@@ -271,6 +375,85 @@ 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,
@@ -279,13 +462,33 @@ 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: false,
+            insertion_debounce: true,
             camera: Default::default(),
             instance_buffer,
             instance_count,
@@ -307,6 +510,19 @@ 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);
     }
@@ -326,6 +542,35 @@ 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 {
@@ -375,6 +620,18 @@ 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)
@@ -382,6 +639,43 @@ 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).
@@ -436,8 +730,88 @@ impl State
             mapped_at_creation: false,
         });
 
-        // ~~ Main render pass ~~
+        // ~~ 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 renderpass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
                 label: None,
@@ -471,15 +845,18 @@ impl State
                 occlusion_query_set: None,
                 multiview_mask: None,
             });
-
-            renderpass.set_pipeline(&self.pipeline);
-            renderpass.set_vertex_buffer(0, self.instance_buffer.slice(..));
+            renderpass.set_vertex_buffer(0, 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));
 
@@ -492,13 +869,10 @@ impl State
         // ~~ EGUI Render pass ~~
         {
             self.egui_renderer.begin_frame(&self.window);
-            egui::Window::new("Window ! ").resizable(true).show(
+            egui::Window::new("Window ! ").resizable(false).show(
                 self.egui_renderer.context(),
                 |ui| {
-                    if self
-                        .camera
-                        .pressed_keyset
-                        .contains(&winit::keyboard::KeyCode::KeyF)
+                    if !self.insertion_debounce
                     {
                         ui.label(
                             egui::RichText::new("Cache paused")
@@ -561,6 +935,23 @@ impl State
             },
         );
 
+        // ~~ Do cache managment
+        // if !self
+        //     .camera
+        //     .pressed_keyset
+        //     .contains(&winit::keyboard::KeyCode::KeyF)
+        if self.insertion_debounce
+        {
+            self.voxel_cache
+                .lock()
+                .cache_post_render(&mut encoder, &self.cache_interface);
+        }
+
+        // ~~ Submit command buffer ~~
+        self.queue.submit([encoder.finish()]);
+        self.window.pre_present_notify();
+        self.queue.present(surface_texture);
+
         // ~~ Get Ray-marching timestamps, report time ~~
         let cloned_rm_time = self.rm_time.clone();
         let cloned_queue = self.queue.clone();
@@ -577,28 +968,12 @@ impl State
         );
 
         // ~~ Do cache managment
-        if !self
-            .camera
-            .pressed_keyset
-            .contains(&winit::keyboard::KeyCode::KeyF)
+        // if !self
+        //     .camera
+        //     .pressed_keyset
+        //     .contains(&winit::keyboard::KeyCode::KeyF)
+        if self.insertion_debounce
         {
-            self.voxel_cache
-                .lock()
-                .cache_post_render(&mut encoder, &self.cache_interface);
-        }
-
-        // ~~ Submit command buffer ~~
-        self.queue.submit([encoder.finish()]);
-        self.window.pre_present_notify();
-        self.queue.present(surface_texture);
-
-        // ~~ Do cache managment
-        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

@@ -3,6 +3,8 @@ 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;
@@ -310,7 +312,7 @@ where
         let heightmap_max = heightmap.iter().copied().reduce(f32::max).unwrap();
 
         heightmap
-            .iter_mut()
+            .par_iter_mut()
             .filter(|x| **x == -9999.)
             .for_each(|x| *x = heightmap_min);
 

src/voxel_cache/structure_table.rs

@@ -80,7 +80,7 @@ impl StructureTable
         &mut self.request_buffer
     }
 
-    pub fn remove_structure(&mut self, structure_id: u32)
+    pub fn free_structure(&mut self, structure_id: u32)
     {
         self.available_slots += 1;
         self.allocation_table[structure_id as usize] = false;
@@ -104,8 +104,7 @@ impl StructureTable
             .allocation_table
             .iter()
             .enumerate()
-            .filter(|(_, allocated)| !**allocated)
-            .next()
+            .find(|(_, allocated)| !**allocated)
             .unwrap();
         self.allocation_table[first_id] = true;
         self.available_slots -= 1;