Works

.gitignore

@@ -1,6 +1,11 @@
 /target
 Cargo.lock
 
+img.jpg
+img.png
+img_low.jpg
+imgs.tar.gz
+vxls_height.tif
 
 # Added by cargo
 #

Cargo.toml

@@ -14,6 +14,7 @@ env_logger = "0.11.11"
 fastapprox = "0.3.1"
 glam = "0.33.5"
 image = "0.25.10"
+indicatif = "0.18.6"
 itertools = "0.15.0"
 pollster = "1.0.1"
 rand = "0.10.2"

shaders/voxel.wgsl

@@ -216,8 +216,9 @@ 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.;
-    let cone_factor = tan((fovy_deg / 180.) * 3.14159) * 2.;
+    let fovy_deg = 100. / 1920.;
+    let fovy_rad = (fovy_deg * 3.14) / 180.;
+    let cone_factor = tan(fovy_rad / 2.) * 2.;
 
     let st_pointer = structure_table_pointer[root_id];
 
@@ -250,25 +251,32 @@ fn new_traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_id: u32, dist_of
 
     // Start location
     //let voxel_dir = select(vec3(-1), vec3(1), ray_dir >= vec3(0.));
-    var node_size = 1 << u32(((max_depth - current_depth) * 2));
-    var child_size = node_size / 4;
+    var node_shift = (max_depth - current_depth) * 2;
+
+    var child_size = 1 << u32(node_shift - 2);
+    var node_size = 1 << u32(node_shift);
+
     var pos_origin = clamp(ray_origin * f32(1 << u32(max_depth * 2)), vec3(0.), vec3(f32(node_size) - 1.));
     var voxel = vec3<i32>(pos_origin);
     var far_t = 0.;
+    var inv_ray_dir = 1. / ray_dir;
+    var ray_positive = ray_dir > vec3(0.);
+    var step_dir = select(vec3(-1), vec3(1), ray_positive);
 
     for(var iter = 0; iter < 400; iter ++)
     {
         // Compute child position
-        node_size = 1 << u32(((max_depth - current_depth) * 2));
-        child_size = node_size / 4;
-        var child_pos = (voxel / child_size) % 4;
-        var pointer = structure_pool[dfs_stack[current_depth]].pointers[child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4];
+        node_shift = (max_depth - current_depth) * 2;
+        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];
+        
+        let min_child_size = (length(vec3<f32>(voxel) - pos_origin) + dist_offset_voxel) * cone_factor;
         while(node_subdivided(pointer) && 
-            !((length(vec3<f32>(voxel) - pos_origin) + dist_offset) * cone_factor >= f32(node_size / 4))
+            f32(child_size / 4) >= min_child_size
             )
         {
-
             if(!node_pointer_valid(pointer) && node_subdivided(pointer))
             {
                 // Record request
@@ -279,9 +287,9 @@ fn new_traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_id: u32, dist_of
             // Descend 
             current_depth += 1;
 
-            node_size /= 4;
-            child_size /= 4;
-            child_pos = (voxel / child_size) % 4;
+            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);
 
             pointer = structure_pool[dfs_stack[current_depth]].pointers[child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4];
@@ -302,13 +310,14 @@ fn new_traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_id: u32, dist_of
         }
 
         // Advance
-        child_pos = (voxel / child_size) * child_size;
-        let far_wall = child_pos + select(vec3(0), vec3(child_size), ray_dir > vec3(0.));
-        let far_wall_inter = (vec3<f32>(far_wall) - pos_origin) / ray_dir;
+        child_pos = voxel & vec3(i32(0xFFFFFFFF << u32(node_shift - 2)));
+        let far_wall = child_pos + select(vec3(0), vec3(child_size), ray_positive);
+        let far_wall_inter = (vec3<f32>(far_wall) - pos_origin) * inv_ray_dir;
         far_t = min(min(far_wall_inter.x, far_wall_inter.y), far_wall_inter.z);
 
         // Perform dda step on the children scale
-        let next_child = select(child_pos, child_pos + select(vec3(-1), vec3(1), ray_dir > vec3(0.)) * vec3(child_size), vec3(far_t) == far_wall_inter);
+        //let next_child = select(child_pos, child_pos + select(vec3(-1), vec3(1), ray_dir > vec3(0.)) * vec3(child_size), vec3(far_t) == far_wall_inter);
+        let next_child = select(child_pos, child_pos + step_dir * vec3(child_size), vec3(far_t) == far_wall_inter);
         
         let previous_voxel = voxel;
         voxel = clamp(vec3<i32>(pos_origin + far_t * ray_dir), next_child, next_child + vec3(child_size) - vec3(1));
@@ -549,10 +558,11 @@ fn fragment(in: VertexOutput) -> FragmentOutput
     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));
+    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;

src/main.rs

@@ -1,4 +1,5 @@
 #![feature(generic_const_exprs)]
+#![feature(float_algebraic)]
 
 use core::sync;
 use std::cell::RefCell;
@@ -12,6 +13,7 @@ use std::sync::mpsc::sync_channel;
 
 use bytemuck::Pod;
 use bytemuck::Zeroable;
+use bytemuck::cast_slice;
 use crevice::std140::AsStd140;
 use crevice::std430::AsStd430;
 use egui::Color32;
@@ -78,26 +80,18 @@ use crate::producers::ChunkedProducer;
 use crate::producers::Producer;
 use crate::producers::SineGenerator;
 use crate::producers::TerrainGenerator;
-use crate::voxel::cache::CacheNodeRequest;
-use crate::voxel::cache::CacheResponse;
-use crate::voxel::cache::ColorBytes;
-use crate::voxel::cache::DestinationElement;
-use crate::voxel::cache::LocationPoolElement;
-use crate::voxel::cache::RequestBuffer;
-use crate::voxel::cache::UsageBuffer;
-use crate::voxel::cache::VoxelCache;
-use crate::voxel::gpu::ExplicitNTreeNode;
-use crate::voxel::pipeline::ChunkHandle;
-use crate::voxel::pipeline::ChunkObject;
-use crate::voxel::pipeline::VoxelPipeline;
-use crate::voxel::sparse::Color;
-use crate::voxel::sparse::NTree;
-use crate::voxel::sparse::NTreeNodeLocator;
+use crate::voxel_cache::cache::CacheNodeRequest;
+use crate::voxel_cache::cache::CacheResponse;
+use crate::voxel_cache::cache::ColorBytes;
+use crate::voxel_cache::cache::DestinationElement;
+use crate::voxel_cache::cache::LocationPoolElement;
+use crate::voxel_cache::cache::VoxelCache;
+use crate::voxel_cache::sparse::NTreeNodeLocator;
 
 mod camera;
 mod egui_renderer;
 mod producers;
-mod voxel;
+mod voxel_cache;
 //mod tree;
 //
 
@@ -120,6 +114,7 @@ struct State
     instance_buffer: Buffer,
     instance_count: usize,
     usage_vec: Arc<Mutex<Vec<usize>>>,
+    rm_time: Arc<Mutex<f32>>,
     insertion_debounce: bool,
 
     camera: Camera,
@@ -167,7 +162,9 @@ impl State
             .unwrap();
         let (device, queue) = adapter
             .request_device(&wgpu::DeviceDescriptor {
-                required_features: Features::IMMEDIATES | Features::SHADER_EARLY_DEPTH_TEST,
+                required_features: Features::IMMEDIATES
+                    | Features::SHADER_EARLY_DEPTH_TEST
+                    | Features::TIMESTAMP_QUERY,
                 required_limits: wgpu::Limits {
                     max_immediate_size: 112,
                     max_storage_buffers_per_shader_stage: 16,
@@ -186,9 +183,21 @@ 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 terrain_generator = TerrainGenerator::<4>::new(5, "vxls_height.tif", 0.2, "img.jpg");
+        //let terrain_generator = TerrainGenerator::<4>::new(5, "vxls_height.tif", 0.2, "img.jpg");
+        // let terrain_generator = TerrainGenerator::<4>::new(
+        //     5,
+        //     "./pointe_percee/height.tif",
+        //     0.2,
+        //     "./pointe_percee/ortho.jpg",
+        // );
+        let terrain_generator = TerrainGenerator::<4>::new(
+            5,
+            "/home/albin/Documents/vxls_maps/lapiz/height.tif",
+            0.2,
+            "/home/albin/Documents/vxls_maps/lapiz/ortho.jpg",
+        );
 
         let mut chunk_pos_map = HashMap::new();
         let chunk_instances = (0..terrain_generator.chunk_width)
@@ -305,6 +314,7 @@ impl State
             instance_count,
             terrain_generator: Arc::new(terrain_generator),
             chunk_pos_map: chunk_pos_map.into(),
+            rm_time: Arc::new(Mutex::new(0.)),
         };
 
         // Configure surface for the first time
@@ -425,6 +435,19 @@ impl State
             }
         };
 
+        let timestamp_query = self.device.create_query_set(&wgpu::QuerySetDescriptor {
+            label: Some("timestamp_query_set"),
+            ty: wgpu::QueryType::Timestamp,
+            count: 2,
+        });
+
+        let timestamp_buffer = self.device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("timestamp_buffer"),
+            size: (size_of::<u64>() * 2) as u64,
+            usage: BufferUsages::QUERY_RESOLVE | BufferUsages::COPY_SRC,
+            mapped_at_creation: false,
+        });
+
         let texture_view = surface_texture
             .texture
             .create_view(&wgpu::TextureViewDescriptor {
@@ -457,7 +480,11 @@ impl State
                     }),
                     stencil_ops: None,
                 }),
-                timestamp_writes: None,
+                timestamp_writes: Some(wgpu::RenderPassTimestampWrites {
+                    query_set: &timestamp_query,
+                    beginning_of_pass_write_index: Some(0),
+                    end_of_pass_write_index: Some(1),
+                }),
                 occlusion_query_set: None,
                 multiview_mask: None,
             });
@@ -475,11 +502,13 @@ impl State
 
             // End the renderpass.
             drop(renderpass);
+
+            encoder.resolve_query_set(&timestamp_query, 0..2, &timestamp_buffer, 0);
         }
 
         let requests = self.device.create_buffer(&wgpu::BufferDescriptor {
             label: Some("dummy_dumb_dinky_aaaahhh_buffer"),
-            size: 16 * 1024,
+            size: 16 * 2000,
             usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
             mapped_at_creation: false,
         });
@@ -511,6 +540,11 @@ impl State
                                 .size(28.),
                         );
                     }
+
+                    ui.label(format!(
+                        "Ray-marching time: {}",
+                        *self.rm_time.lock() / 1_000_000.
+                    ));
                     egui_plot::Plot::new("Plot").show(ui, |plot_ui| {
                         plot_ui.bar_chart(BarChart::new(
                             "histo",
@@ -580,6 +614,22 @@ impl State
         //     .contains(&winit::keyboard::KeyCode::KeyF))
         //     && !self.insertion_debounce
         // {
+
+        // Report frame time
+        let cloned_rm_time = self.rm_time.clone();
+        let cloned_queue = self.queue.clone();
+        DownloadBuffer::read_buffer(
+            &self.device,
+            &self.queue,
+            &timestamp_buffer.slice(..),
+            move |buffer| {
+                let buffer_slice = buffer.unwrap();
+                let slice: &[u64] = cast_slice(&buffer_slice);
+                let time = (slice[1] - slice[0]) as f32 * cloned_queue.get_timestamp_period();
+                *cloned_rm_time.lock() = time;
+            },
+        );
+
         if !self
             .camera
             .pressed_keyset

src/producers.rs

@@ -2,11 +2,12 @@ use std::fs::File;
 use std::path::Path;
 
 use glam::Vec3;
+use indicatif::ProgressIterator;
 use itertools::Itertools;
 
-use crate::voxel::gpu::ExplicitNTreeNode;
-use crate::voxel::gpu::StructurePointer;
-use crate::voxel::sparse::Color;
+use crate::voxel_cache::gpu::ExplicitNTreeNode;
+use crate::voxel_cache::gpu::StructurePointer;
+use crate::voxel_cache::sparse::Color;
 
 pub struct BallGenerator<const N: usize>
 {
@@ -15,7 +16,13 @@ pub struct BallGenerator<const N: usize>
 
 pub fn map(x: f32, x_min: f32, x_max: f32, y_min: f32, y_max: f32) -> f32
 {
-    ((x - x_min) / (x_max - x_min)) * (y_max - y_min) + y_min
+    //((x - x_min) / (x_max - x_min)) * (y_max - y_min) + y_min
+    let input_range = x_max.algebraic_sub(x_min);
+    let output_range = y_max.algebraic_sub(y_min);
+
+    (x.algebraic_sub(x_min).algebraic_div(input_range))
+        .algebraic_mul(output_range)
+        .algebraic_add(y_min)
 }
 
 pub trait Producer<const N: usize>
@@ -240,6 +247,14 @@ where
     heightmap: Vec<f32>,
     colormap: Vec<u8>,
 
+    heightmap_low_width: usize,
+    heightmap_low_height: usize,
+    heightmap_low: Vec<(f32, f32)>,
+
+    colormap_low_width: usize,
+    colormap_low_height: usize,
+    colormap_mip: Vec<u8>,
+
     pub chunk_width: usize,
     pub chunk_height: usize,
     pub chunk_alt: usize,
@@ -256,29 +271,50 @@ where
         color_path: P,
     ) -> Self
     {
+        println!("Starting terrain producer");
+        println!("Loading height map.");
         let mut tiff_dec = tiff::decoder::Decoder::new(File::open(height_path).unwrap()).unwrap();
         let (heightmap_width, heightmap_height) = tiff_dec.dimensions().unwrap();
         let (heightmap_width, heightmap_height) =
             (heightmap_width as usize, heightmap_height as usize);
 
-        let heightmap = match tiff_dec.read_image().unwrap()
+        let mut heightmap = match tiff_dec.read_image().unwrap()
         {
             tiff::decoder::DecodingResult::F32(vec) => vec,
             _ => panic!("Unsupported format"),
         };
 
+        println!("Loading color map.");
         let mut color = image::ImageReader::open(color_path).unwrap();
         color.no_limits();
 
         let color = color.decode().unwrap();
-        let colormap = color.as_rgb8().unwrap().to_vec();
+        let mut colormap = color.as_rgb8().unwrap().to_vec();
+
+        println!("Converting color spaces");
+        colormap.iter_mut().for_each(|x| {
+            let normalized = map(*x as f32, 0., 255., 0., 1.);
+            let maped = normalized.powf(2.4);
+            *x = map(maped, 0., 1., 0., 255.) as u8;
+        });
 
         let terrain_width = color.width() as usize;
         let terrain_height = color.height() as usize;
 
-        let heightmap_min = heightmap.iter().copied().reduce(f32::min).unwrap();
+        println!("Computing heightmap min/max");
+        let heightmap_min = heightmap
+            .iter()
+            .copied()
+            .filter(|x| *x != -9999.)
+            .reduce(f32::min)
+            .unwrap();
         let heightmap_max = heightmap.iter().copied().reduce(f32::max).unwrap();
 
+        heightmap
+            .iter_mut()
+            .filter(|x| **x == -9999.)
+            .for_each(|x| *x = heightmap_min);
+
         // Decide size in chunks
         let height_amplitude = heightmap_max - heightmap_min;
         let chunk_size = N.pow(chunk_power as u32);
@@ -286,6 +322,65 @@ where
         let chunk_height = terrain_height.div_ceil(chunk_size);
         let chunk_alt = ((height_amplitude / height_factor) as usize).div_ceil(chunk_size);
 
+        // build the low heightmap
+        println!("Computing low res height/color maps");
+        let heightmap_low_width = heightmap_width / 8;
+        let heightmap_low_height = heightmap_height / 8;
+
+        let mut heightmap_low = vec![(0., 0.); heightmap_low_height * heightmap_low_width];
+
+        for y in 0..heightmap_low_height
+        {
+            for x in 0..heightmap_low_width
+            {
+                let mut min = heightmap_max;
+                let mut max = heightmap_min;
+                for sy in (y * 8)..(y * 8 + 8)
+                {
+                    for sx in (x * 8)..(x * 8 + 8)
+                    {
+                        min = min.min(heightmap[sx + sy * heightmap_width]);
+                        max = max.max(heightmap[sx + sy * heightmap_width]);
+                    }
+                }
+
+                heightmap_low[x + y * heightmap_low_width] = (min, max);
+            }
+        }
+
+        // build the color map mip
+        let colormap_low_width = terrain_width / 8;
+        let colormap_low_height = terrain_height / 8;
+
+        let mut colormap_mip = vec![0u8; colormap_low_width * colormap_low_height * 3];
+
+        for y in 0..colormap_low_height
+        {
+            for x in 0..colormap_low_width
+            {
+                let mut r = 0u32;
+                let mut g = 0u32;
+                let mut b = 0u32;
+                for sy in (y * 8)..(y * 8 + 8)
+                {
+                    for sx in (x * 8)..(x * 8 + 8)
+                    {
+                        r += colormap[(sx + sy * terrain_width) * 3] as u32;
+                        g += colormap[(sx + sy * terrain_width) * 3 + 1] as u32;
+                        b += colormap[(sx + sy * terrain_width) * 3 + 2] as u32;
+                    }
+                }
+
+                colormap_mip[(x + y * colormap_low_width) * 3] =
+                    ((r as f32) / (8 * 8) as f32).clamp(0., 255.) as u8;
+                colormap_mip[(x + y * colormap_low_width) * 3 + 1] =
+                    ((g as f32) / (8 * 8) as f32).clamp(0., 255.) as u8;
+                colormap_mip[(x + y * colormap_low_width) * 3 + 2] =
+                    ((b as f32) / (8 * 8) as f32).clamp(0., 255.) as u8;
+            }
+        }
+
+        println!("Producer ready");
         Self {
             chunk_power,
             heightmap_width,
@@ -297,6 +392,14 @@ where
             heightmap,
             colormap,
 
+            heightmap_low_width,
+            heightmap_low_height,
+            heightmap_low,
+
+            colormap_low_width,
+            colormap_low_height,
+            colormap_mip,
+
             chunk_width,
             chunk_height,
             chunk_alt,
@@ -340,30 +443,68 @@ where
             let mut sample_min = self.heightmap_max;
             let mut color_avg = Color(0., 0., 0., 0.);
             let mut count = 0;
-            for (x, z) in (0..child_size).cartesian_product(0..child_size)
+
+            if depth <= 2
             {
-                let gvx = gcx + x;
-                let gvz = gcz + z;
-                if gvx < self.terrain_width && gvz < self.terrain_height
+                for (z, x) in (0..(child_size / 8)).cartesian_product(0..(child_size / 8))
                 {
-                    // Height sample
-                    let height_x = (gvx * self.heightmap_width) / self.terrain_width;
-                    let height_z = (gvz * self.heightmap_height) / self.terrain_height;
+                    let gvx = (gcx + x * 8) / 8;
+                    let gvz = (gcz + z * 8) / 8;
+                    if gvx < self.colormap_low_width && gvz < self.colormap_low_height
+                    {
+                        // Height sample
+                        let height_x = (gvx * self.heightmap_low_width) / self.colormap_low_width;
+                        let height_z = (gvz * self.heightmap_low_height) / self.colormap_low_height;
 
-                    let sample = self.heightmap[height_x + height_z * self.heightmap_width];
-                    sample_max = sample_max.max(sample);
-                    sample_min = sample_min.min(sample);
+                        let sample =
+                            self.heightmap_low[height_x + height_z * self.heightmap_low_width];
+                        sample_min = sample_min.min(sample.0);
+                        sample_max = sample_max.max(sample.1);
 
-                    let sample_color_r = self.colormap[(gvx + gvz * self.terrain_width) * 3];
-                    let sample_color_g = self.colormap[(gvx + gvz * self.terrain_width) * 3 + 1];
-                    let sample_color_b = self.colormap[(gvx + gvz * self.terrain_width) * 3 + 2];
+                        let sample_color_r =
+                            self.colormap_mip[(gvx + gvz * self.colormap_low_width) * 3];
+                        let sample_color_g =
+                            self.colormap_mip[(gvx + gvz * self.colormap_low_width) * 3 + 1];
+                        let sample_color_b =
+                            self.colormap_mip[(gvx + gvz * self.colormap_low_width) * 3 + 2];
 
-                    color_avg.0 += map(sample_color_r as f32, 0., 256., 0., 1.);
-                    color_avg.1 += map(sample_color_g as f32, 0., 256., 0., 1.);
-                    color_avg.2 += map(sample_color_b as f32, 0., 256., 0., 1.);
-                    count += 1;
+                        color_avg.0 += map(sample_color_r as f32, 0., 256., 0., 1.);
+                        color_avg.1 += map(sample_color_g as f32, 0., 256., 0., 1.);
+                        color_avg.2 += map(sample_color_b as f32, 0., 256., 0., 1.);
+                        count += 1;
+                    }
+                    //let gvy = gcy;
+                }
+            }
+            else
+            {
+                for (z, x) in (0..child_size).cartesian_product(0..child_size)
+                {
+                    let gvx = gcx + x;
+                    let gvz = gcz + z;
+                    if gvx < self.terrain_width && gvz < self.terrain_height
+                    {
+                        // Height sample
+                        let height_x = (gvx * self.heightmap_width) / self.terrain_width;
+                        let height_z = (gvz * self.heightmap_height) / self.terrain_height;
+
+                        let sample = self.heightmap[height_x + height_z * self.heightmap_width];
+                        sample_max = sample_max.max(sample);
+                        sample_min = sample_min.min(sample);
+
+                        let sample_color_r = self.colormap[(gvx + gvz * self.terrain_width) * 3];
+                        let sample_color_g =
+                            self.colormap[(gvx + gvz * self.terrain_width) * 3 + 1];
+                        let sample_color_b =
+                            self.colormap[(gvx + gvz * self.terrain_width) * 3 + 2];
+
+                        color_avg.0 += map(sample_color_r as f32, 0., 256., 0., 1.);
+                        color_avg.1 += map(sample_color_g as f32, 0., 256., 0., 1.);
+                        color_avg.2 += map(sample_color_b as f32, 0., 256., 0., 1.);
+                        count += 1;
+                    }
+                    //let gvy = gcy;
                 }
-                //let gvy = gcy;
             }
 
             color_avg.0 /= count as f32;

src/voxel/pipeline.rs

@@ -1,799 +0,0 @@
-use std::num::NonZero;
-
-use bytemuck::cast_slice;
-use crevice::std140::AsStd140;
-use glam::Mat4;
-use wgpu::BindGroup;
-use wgpu::BindGroupDescriptor;
-use wgpu::BindGroupEntry;
-use wgpu::BindGroupLayout;
-use wgpu::Buffer;
-use wgpu::BufferUsages;
-use wgpu::CommandEncoder;
-use wgpu::CommandEncoderDescriptor;
-use wgpu::ComputePass;
-use wgpu::ComputePassDescriptor;
-use wgpu::ComputePipeline;
-use wgpu::Device;
-use wgpu::Operations;
-use wgpu::Queue;
-use wgpu::RenderPipeline;
-use wgpu::ShaderModuleDescriptor;
-use wgpu::ShaderStages;
-use wgpu::TextureFormat;
-use wgpu::TextureView;
-use wgpu::VertexBufferLayout;
-use wgpu::util::StagingBelt;
-
-use crate::as_raw_bytes;
-use crate::camera::Camera;
-use crate::voxel::cache::ColorPoolElement;
-use crate::voxel::cache::LocationPoolElement;
-use crate::voxel::cache::RequestBufferElement;
-use crate::voxel::cache::StructurePoolElement;
-use crate::voxel::gpu::StructurePointer;
-use crate::voxel::sparse::Color;
-
-// Represents a chunk to be rendered by the voxel pipeline
-#[derive(Clone, Copy)]
-#[repr(C)]
-pub struct ChunkObject
-{
-    // Chunk object transform
-    pub transform: Mat4,
-
-    // Chunk data
-    pub color: Color,
-    pub subdivided: bool,
-
-    // Producer specific data
-    pub id: u32,
-}
-
-#[derive(Clone, Copy)]
-#[repr(C)]
-pub struct CacheChunkObject
-{
-    // Chunk object transform
-    transform: Mat4,
-
-    // Chunk data
-    color: Color,
-
-    // Producer specific data
-    id: u32,
-
-    // Pointer into cache
-    pointer: StructurePointer,
-}
-
-#[derive(Clone, Copy)]
-pub struct ChunkHandle(usize);
-
-pub struct CacheRequest
-{
-    // Records how many rays requested a
-    // resource
-    count: u32,
-}
-
-pub struct VoxelPipeline<const N: usize>
-where
-    [(); N * N * N]:,
-{
-    chunk_allocations: Vec<bool>,
-    chunk_indices: Buffer,
-    chunk_staging: StagingBelt,
-    chunk_objects: Buffer,
-    chunk_requests: Buffer,
-
-    chunk_objects_bind_group_layout: BindGroupLayout,
-    ray_bind_group_layout: BindGroupLayout,
-    chunk_objects_bind_group: BindGroup,
-    chunk_requests_bind_group: BindGroup,
-    render_pipeline: RenderPipeline,
-
-    // Cache pools
-    structure_pool: Buffer,
-    color_pool: Buffer,
-    location_pool: Buffer,
-
-    // Cache interaction
-    request_buffer: Buffer,
-    usage_buffer: Buffer,
-    cache_bind_group: BindGroup,
-
-    device: Device,
-    queue: Queue,
-
-    // Cache keeping shaders
-    clear_chunk_request: ComputePipeline,
-    sort_requests: ComputePipeline,
-}
-
-#[derive(AsStd140)]
-struct RenderPipelineImmediate
-{
-    view_proj: Mat4,
-}
-
-impl<const N: usize> VoxelPipeline<N>
-where
-    [(); N * N * N]:,
-{
-    pub fn new(
-        cache_size: usize,
-        device: Device,
-        queue: Queue,
-        surface_format: TextureFormat,
-    ) -> Self
-    {
-        let shader_module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
-            label: Some("Main shader module"),
-            source: wgpu::ShaderSource::Wgsl(
-                std::fs::read_to_string("shaders/voxel.wgsl")
-                    .unwrap()
-                    .into(),
-            ),
-        });
-
-        let cache_bind_group_layout =
-            device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
-                label: Some("Cache bind group layout"),
-                entries: &[
-                    // Location pool
-                    wgpu::BindGroupLayoutEntry {
-                        binding: 0,
-                        visibility: ShaderStages::FRAGMENT,
-                        ty: wgpu::BindingType::Buffer {
-                            ty: wgpu::BufferBindingType::Storage { read_only: false },
-                            has_dynamic_offset: false,
-                            min_binding_size: None,
-                        },
-                        count: None,
-                    },
-                    // Color pool
-                    wgpu::BindGroupLayoutEntry {
-                        binding: 1,
-                        visibility: ShaderStages::FRAGMENT,
-                        ty: wgpu::BindingType::Buffer {
-                            ty: wgpu::BufferBindingType::Storage { read_only: false },
-                            has_dynamic_offset: false,
-                            min_binding_size: None,
-                        },
-                        count: None,
-                    },
-                    // Location pool
-                    wgpu::BindGroupLayoutEntry {
-                        binding: 2,
-                        visibility: ShaderStages::FRAGMENT,
-                        ty: wgpu::BindingType::Buffer {
-                            ty: wgpu::BufferBindingType::Storage { read_only: false },
-                            has_dynamic_offset: false,
-                            min_binding_size: None,
-                        },
-                        count: None,
-                    },
-                    // Request buffer
-                    wgpu::BindGroupLayoutEntry {
-                        binding: 3,
-                        visibility: ShaderStages::FRAGMENT,
-                        ty: wgpu::BindingType::Buffer {
-                            ty: wgpu::BufferBindingType::Storage { read_only: false },
-                            has_dynamic_offset: false,
-                            min_binding_size: None,
-                        },
-                        count: None,
-                    },
-                    // Usage buffer
-                    wgpu::BindGroupLayoutEntry {
-                        binding: 4,
-                        visibility: ShaderStages::FRAGMENT,
-                        ty: wgpu::BindingType::Buffer {
-                            ty: wgpu::BufferBindingType::Storage { read_only: false },
-                            has_dynamic_offset: false,
-                            min_binding_size: None,
-                        },
-                        count: None,
-                    },
-                ],
-            });
-
-        let chunk_objects_bind_group_layout =
-            device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
-                label: Some("Chunk objects bg"),
-                entries: &[wgpu::BindGroupLayoutEntry {
-                    binding: 0,
-                    visibility: ShaderStages::VERTEX,
-                    ty: wgpu::BindingType::Buffer {
-                        ty: wgpu::BufferBindingType::Storage { read_only: true },
-                        has_dynamic_offset: false,
-                        min_binding_size: None,
-                    },
-                    count: None,
-                }],
-            });
-
-        let chunk_requests_bind_group_layout =
-            device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
-                label: Some("Ray bind group layout"),
-                entries: &[wgpu::BindGroupLayoutEntry {
-                    binding: 0,
-                    visibility: ShaderStages::FRAGMENT | ShaderStages::COMPUTE,
-                    ty: wgpu::BindingType::Buffer {
-                        ty: wgpu::BufferBindingType::Storage { read_only: false },
-                        has_dynamic_offset: false,
-                        min_binding_size: None,
-                    },
-                    count: None,
-                }],
-            });
-
-        let request_buffer_sort_bind_group_layout =
-            device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
-                label: Some("Ray bind group layout"),
-                entries: &[wgpu::BindGroupLayoutEntry {
-                    binding: 0,
-                    visibility: ShaderStages::COMPUTE,
-                    ty: wgpu::BindingType::Buffer {
-                        ty: wgpu::BufferBindingType::Storage { read_only: false },
-                        has_dynamic_offset: false,
-                        min_binding_size: None,
-                    },
-                    count: None,
-                }],
-            });
-
-        let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
-            label: Some("Voxel pipeline layout"),
-
-            bind_group_layouts: &[
-                Some(&chunk_objects_bind_group_layout),
-                Some(&chunk_requests_bind_group_layout),
-                Some(&cache_bind_group_layout),
-            ],
-            immediate_size: RenderPipelineImmediate::std140_size_static() as u32,
-        });
-
-        let chunk_pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
-            label: Some("Render pipeline"),
-            layout: Some(&pipeline_layout),
-            vertex: wgpu::VertexState {
-                module: &shader_module,
-                entry_point: Some("chunk"),
-                compilation_options: Default::default(),
-                buffers: &[Some(VertexBufferLayout {
-                    array_stride: size_of::<u32>() as u64,
-                    step_mode: wgpu::VertexStepMode::Instance,
-                    attributes: &[wgpu::VertexAttribute {
-                        format: wgpu::VertexFormat::Uint32,
-                        offset: 0,
-                        shader_location: 0,
-                    }],
-                })],
-            },
-            primitive: wgpu::PrimitiveState {
-                topology: wgpu::PrimitiveTopology::TriangleList,
-                strip_index_format: None,
-                front_face: wgpu::FrontFace::Ccw,
-                cull_mode: None,
-                unclipped_depth: false,
-                polygon_mode: wgpu::PolygonMode::Fill,
-                conservative: false,
-            },
-            depth_stencil: Some(wgpu::DepthStencilState {
-                format: wgpu::TextureFormat::Depth24PlusStencil8,
-                depth_write_enabled: Some(true),
-                depth_compare: Some(wgpu::CompareFunction::LessEqual),
-                stencil: wgpu::StencilState::default(),
-                bias: wgpu::DepthBiasState::default(),
-            }),
-            multisample: wgpu::MultisampleState::default(),
-            fragment: Some(wgpu::FragmentState {
-                module: &shader_module,
-                entry_point: Some("fragment"),
-                compilation_options: wgpu::PipelineCompilationOptions::default(),
-                targets: &[Some(wgpu::ColorTargetState {
-                    format: surface_format,
-                    blend: None,
-                    write_mask: wgpu::ColorWrites::default(),
-                })],
-            }),
-            multiview_mask: None,
-            cache: None,
-        });
-
-        let chunk_indices = device.create_buffer(&wgpu::wgt::BufferDescriptor {
-            label: Some("Chunk index buffer"),
-            size: size_of::<u32>() as u64,
-            usage: BufferUsages::VERTEX | BufferUsages::COPY_DST,
-            mapped_at_creation: true,
-        });
-        chunk_indices.unmap();
-
-        let chunk_objects = device.create_buffer(&wgpu::wgt::BufferDescriptor {
-            label: Some("Chunk buffer"),
-            size: size_of::<CacheChunkObject>() as u64,
-            usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC,
-            mapped_at_creation: false,
-        });
-
-        let chunk_requests = device.create_buffer(&wgpu::wgt::BufferDescriptor {
-            label: Some("Chunk request buffer"),
-            size: size_of::<u32>() as u64,
-            usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC,
-            mapped_at_creation: false,
-        });
-
-        // Pools
-        let structure_pool = device.create_buffer(&wgpu::wgt::BufferDescriptor {
-            label: Some("Structure pool"),
-            size: size_of::<StructurePoolElement<N>>() as u64 * cache_size as u64,
-            usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC,
-            mapped_at_creation: false,
-        });
-
-        let color_pool = device.create_buffer(&wgpu::wgt::BufferDescriptor {
-            label: Some("Color pool"),
-            size: size_of::<ColorPoolElement<N>>() as u64 * cache_size as u64,
-            usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC,
-            mapped_at_creation: false,
-        });
-
-        let location_pool = device.create_buffer(&wgpu::wgt::BufferDescriptor {
-            label: Some("Locatino pool"),
-            size: size_of::<LocationPoolElement>() as u64 * cache_size as u64,
-            usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC,
-            mapped_at_creation: false,
-        });
-
-        let request_buffer = device.create_buffer(&wgpu::wgt::BufferDescriptor {
-            label: Some("Request buffer"),
-            size: size_of::<RequestBufferElement<N>>() as u64 * cache_size as u64,
-            usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC,
-            mapped_at_creation: false,
-        });
-
-        let request_sort_buffer = device.create_buffer(&wgpu::wgt::BufferDescriptor {
-            label: Some("Request sort buffer"),
-            size: size_of::<u32>() as u64 * cache_size as u64,
-            usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC,
-            mapped_at_creation: true,
-        });
-
-        request_sort_buffer
-            .get_mapped_range_mut(0..)
-            .unwrap()
-            .copy_from_slice(cast_slice(
-                (0..cache_size as u32).collect::<Vec<_>>().as_slice(),
-            ));
-        request_sort_buffer.unmap();
-
-        let usage_buffer = device.create_buffer(&wgpu::wgt::BufferDescriptor {
-            label: Some("Usage buffer buffer"),
-            size: size_of::<u32>() as u64 * cache_size as u64,
-            usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC,
-            mapped_at_creation: false,
-        });
-
-        let cache_bind_group = device.create_bind_group(&BindGroupDescriptor {
-            label: Some("Cache bind group"),
-            layout: &cache_bind_group_layout,
-            entries: &[
-                // Structure pool
-                wgpu::BindGroupEntry {
-                    binding: 0,
-                    resource: structure_pool.as_entire_binding(),
-                },
-                // Color pool
-                wgpu::BindGroupEntry {
-                    binding: 1,
-                    resource: color_pool.as_entire_binding(),
-                },
-                // Location pool
-                wgpu::BindGroupEntry {
-                    binding: 2,
-                    resource: location_pool.as_entire_binding(),
-                },
-                // Request buffer
-                wgpu::BindGroupEntry {
-                    binding: 3,
-                    resource: request_buffer.as_entire_binding(),
-                },
-                // Usage buffer
-                wgpu::BindGroupEntry {
-                    binding: 4,
-                    resource: structure_pool.as_entire_binding(),
-                },
-            ],
-        });
-
-        // Cache keeping shaders
-        let clear_chunk_request =
-            device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
-                label: Some("Clean chunk request"),
-                layout: Some(
-                    &device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
-                        label: Some("clear chunk requests layout"),
-                        bind_group_layouts: &[Some(&chunk_requests_bind_group_layout)],
-                        immediate_size: 0,
-                    }),
-                ),
-                module: &device.create_shader_module(ShaderModuleDescriptor {
-                    label: Some("clear_chunk_requests shader module"),
-                    source: wgpu::ShaderSource::Wgsl(
-                        "
-                            @group(0) @binding(0) var<storage, read_write> chunk_requests: array<u32>; 
-                            @compute
-                            @workgroup_size(16)
-                            fn main(
-                                    @builtin(global_invocation_id) global_invocation_id: vec3<u32>
-                                )
-                            {
-                                    let index = global_invocation_id.x;
-                                    let total = arrayLength(&chunk_requests);
-                            
-                                    if(index < total)
-                                    {
-                                        chunk_requests[index] = 0;
-                                    }
-                            }
-                        "
-                        .into(),
-                    ),
-                }),
-                entry_point: Some("main"),
-                compilation_options: Default::default(),
-                cache: None,
-            });
-
-        // Sorting requests
-        let sort_requests_bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
-            label: Some("Sort request bing group"),
-            layout: &request_buffer_sort_bind_group_layout,
-            entries: &[BindGroupEntry {
-                binding: 0,
-                resource: request_sort_buffer.as_entire_binding(),
-            }],
-        });
-        let sort_requests =
-            device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
-                label: Some("Sort node request"),
-                layout: Some(
-                    &device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
-                        label: Some("Sort node requests"),
-                        bind_group_layouts: &[Some(&cache_bind_group_layout), Some(&request_buffer_sort_bind_group_layout)],
-                        immediate_size: 0,
-                    }),
-                ),
-                module: &device.create_shader_module(ShaderModuleDescriptor {
-                    label: Some("clear_chunk_requests shader module"),
-                    source: wgpu::ShaderSource::Wgsl(
-                        "
-                            struct RequestElement
-                            {
-                                children: array<atomic<u32>, 64> 
-                            }
-
-                            @group(0) @binding(0) var<storage, read_write> structure_pool: array<u32>;
-                            @group(0) @binding(1) var<storage, read_write> color_pool: array<u32>;
-                            @group(0) @binding(2) var<storage, read_write> location_pool: array<u32>;
-                            @group(0) @binding(3) var<storage, read_write> request_buffer: array<RequestElement>;
-                            @group(0) @binding(4) var<storage, read_write> usage_buffer: array<u32>;
-
-                            @group(1) @binding(0) var<storage, read_write> sort_indirection: array<u32>; 
-
-                            @compute
-                            @workgroup_size(16)
-                            fn main(
-                                    @builtin(global_invocation_id) global_invocation_id: vec3<u32>
-                                )
-                            {
-                                    let index = global_invocation_id.x;
-                                    let total = arrayLength(&chunk_requests);
-
-                                    // Odd pass
-                                    let a = index * 2 + 1;
-                                    let b = a + 1;
-
-                                    // Gather elements
-                                    let av = request_buffer[sort_indirection[a]];
-                                    let bv = request_buffer[sort_indirection[b]];
-
-                                    if b < total && av > bv
-                                    {
-                                        request_buffer[sort_indirection[a]] = bv;
-                                        request_buffer[sort_indirection[b]] = av;
-                                    }
-                                    storageBarrier();
-
-                                    // Even pass
-                                    a = index * 2;
-                                    b = a + 1;
-
-                                    // Gather elements
-                                    av = request_buffer[sort_indirection[a]];
-                                    bv = request_buffer[sort_indirection[b]];
-
-                                    if b < total && av > bv
-                                    {
-                                        request_buffer[sort_indirection[a]] = bv;
-                                        request_buffer[sort_indirection[b]] = av;
-                                    } 
-                            }
-                        "
-                        .into(),
-                    ),
-                }),
-                entry_point: Some("main"),
-                compilation_options: Default::default(),
-                cache: None,
-            });
-
-        VoxelPipeline {
-            // Only one slot, no allocated chunks at the beginning
-            chunk_allocations: vec![false],
-            chunk_objects_bind_group: device.create_bind_group(&wgpu::BindGroupDescriptor {
-                label: Some("Chunk objects bind group"),
-                layout: &chunk_objects_bind_group_layout,
-                entries: &[wgpu::BindGroupEntry {
-                    binding: 0,
-                    resource: chunk_objects.as_entire_binding(),
-                }],
-            }),
-            chunk_requests_bind_group: device.create_bind_group(&wgpu::BindGroupDescriptor {
-                label: Some("Ray bind group"),
-                layout: &chunk_requests_bind_group_layout,
-                entries: &[wgpu::BindGroupEntry {
-                    binding: 0,
-                    resource: chunk_requests.as_entire_binding(),
-                }],
-            }),
-            chunk_objects,
-            chunk_requests,
-            chunk_indices,
-
-            structure_pool,
-            color_pool,
-            location_pool,
-            request_buffer,
-            usage_buffer,
-            cache_bind_group,
-
-            chunk_objects_bind_group_layout,
-            ray_bind_group_layout: chunk_requests_bind_group_layout,
-            chunk_staging: StagingBelt::new(device.clone(), size_of::<CacheChunkObject>() as u64),
-            render_pipeline: chunk_pipeline,
-            device,
-            queue,
-
-            clear_chunk_request,
-            sort_requests,
-        }
-    }
-
-    pub fn render(
-        &mut self,
-        encoder: &mut CommandEncoder,
-        texture_view: &TextureView,
-        depth_buffer_view: &TextureView,
-        camera: &Camera,
-    )
-    {
-        let mut renderpass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
-            label: None,
-            color_attachments: &[Some(wgpu::RenderPassColorAttachment {
-                view: texture_view,
-                depth_slice: None,
-                resolve_target: None,
-                ops: wgpu::Operations {
-                    load: wgpu::LoadOp::Clear(wgpu::Color::BLACK),
-                    store: wgpu::StoreOp::Store,
-                },
-            })],
-            depth_stencil_attachment: Some(wgpu::RenderPassDepthStencilAttachment {
-                view: depth_buffer_view,
-                depth_ops: Some(Operations {
-                    load: wgpu::LoadOp::Clear(1.),
-                    store: wgpu::StoreOp::Discard,
-                }),
-                stencil_ops: None,
-            }),
-            timestamp_writes: None,
-            occlusion_query_set: None,
-            multiview_mask: None,
-        });
-
-        renderpass.set_pipeline(&self.render_pipeline);
-        renderpass.set_bind_group(0, Some(&self.chunk_objects_bind_group), &[]);
-        renderpass.set_bind_group(1, Some(&self.chunk_requests_bind_group), &[]);
-        renderpass.set_bind_group(2, Some(&self.cache_bind_group), &[]);
-        renderpass.set_vertex_buffer(0, self.chunk_indices.slice(0..));
-        renderpass.set_immediates(
-            0,
-            RenderPipelineImmediate {
-                view_proj: camera.view_proj(),
-            }
-            .as_std140()
-            .as_bytes(),
-        );
-        renderpass.draw(
-            0..36,
-            0..(self.chunk_allocations.iter().filter(|x| **x).count() as u32),
-        );
-
-        // End the renderpass.
-        drop(renderpass);
-
-        let mut compute_pass = encoder.begin_compute_pass(&ComputePassDescriptor {
-            label: Some("cache keeping pass"),
-            timestamp_writes: None,
-        });
-
-        compute_pass.set_bind_group(0, Some(&self.chunk_requests_bind_group), &[]);
-        compute_pass.set_pipeline(&self.clear_chunk_request);
-        compute_pass.dispatch_workgroups(
-            self.chunk_allocations.len().next_multiple_of(16) as u32 / 16,
-            1,
-            1,
-        );
-
-        drop(compute_pass)
-    }
-
-    fn update_indices(&mut self)
-    {
-        let indices = self
-            .chunk_allocations
-            .iter()
-            .enumerate()
-            .filter(|(_, b)| **b)
-            .map(|(i, _)| i as u32)
-            .collect::<Vec<_>>();
-
-        self.chunk_indices = self.device.create_buffer(&wgpu::wgt::BufferDescriptor {
-            label: Some("Chunk index buffer"),
-            size: size_of::<u32>() as u64 * indices.len() as u64,
-            usage: BufferUsages::VERTEX | BufferUsages::COPY_DST,
-            mapped_at_creation: true,
-        });
-
-        // Copy chunk indices into new buffer
-        self.chunk_indices
-            .get_mapped_range_mut(0..(size_of::<u32>() as u64 * indices.len() as u64))
-            .unwrap()
-            .copy_from_slice(cast_slice(&indices));
-        self.chunk_indices.unmap();
-    }
-
-    pub fn remove_chunks(&mut self, handles: &[ChunkHandle])
-    {
-        for handle in handles.iter()
-        {
-            self.chunk_allocations[handle.0] = false;
-        }
-        self.update_indices();
-    }
-
-    pub fn push_new_chunks(&mut self, objects: &[ChunkObject]) -> Vec<ChunkHandle>
-    {
-        // Find room for new chunks
-        let mut destinations = vec![0; objects.len()];
-
-        let mut encoder = self
-            .device
-            .create_command_encoder(&CommandEncoderDescriptor {
-                label: Some("Chunk buffer writes"),
-            });
-
-        // count available space
-        let space = self.chunk_allocations.iter().filter(|x| !*x).count();
-        if space < objects.len()
-        {
-            // Allocate more space
-            // Get first bigger power of two
-            let necessary_space =
-                (objects.len() + self.chunk_allocations.len()).next_power_of_two();
-            dbg!(necessary_space);
-            self.chunk_allocations
-                .extend(vec![false; necessary_space - self.chunk_allocations.len()]);
-
-            // Make buffer bigger
-            let old_buffer = self.chunk_objects.clone();
-            self.chunk_objects = self.device.create_buffer(&wgpu::wgt::BufferDescriptor {
-                label: Some("Chunk buffer"),
-                size: size_of::<ChunkObject>() as u64 * self.chunk_allocations.len() as u64,
-                usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC,
-                mapped_at_creation: false,
-            });
-
-            let old_chunk_requests = self.chunk_requests.clone();
-            self.chunk_requests = self.device.create_buffer(&wgpu::wgt::BufferDescriptor {
-                label: Some("Chunk request buffer"),
-                size: size_of::<u32>() as u64 * self.chunk_allocations.len() as u64,
-                usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC,
-                mapped_at_creation: false,
-            });
-
-            self.chunk_objects_bind_group =
-                self.device.create_bind_group(&wgpu::BindGroupDescriptor {
-                    label: Some("Chunk objects bind group"),
-                    layout: &self.chunk_objects_bind_group_layout,
-                    entries: &[wgpu::BindGroupEntry {
-                        binding: 0,
-                        resource: self.chunk_objects.as_entire_binding(),
-                    }],
-                });
-
-            self.chunk_requests_bind_group =
-                self.device.create_bind_group(&wgpu::BindGroupDescriptor {
-                    label: Some("Ray bind group"),
-                    layout: &self.ray_bind_group_layout,
-                    entries: &[wgpu::BindGroupEntry {
-                        binding: 0,
-                        resource: self.chunk_requests.as_entire_binding(),
-                    }],
-                });
-
-            encoder.copy_buffer_to_buffer(
-                &old_chunk_requests,
-                0,
-                &self.chunk_requests,
-                0,
-                old_chunk_requests.size(),
-            );
-
-            encoder.copy_buffer_to_buffer(
-                &old_buffer,
-                0,
-                &self.chunk_objects,
-                0,
-                old_buffer.size(),
-            );
-        }
-
-        let mut dest_ptr = 0;
-        for (i, allocated) in self
-            .chunk_allocations
-            .iter_mut()
-            .enumerate()
-            .filter(|(_, allocated)| !**allocated)
-            .take(destinations.len())
-        {
-            if !*allocated
-            {
-                *allocated = true;
-                destinations[dest_ptr] = i;
-                dest_ptr += 1;
-            }
-        }
-
-        // Write each new chunk
-        for (destination, object) in destinations.iter().zip(objects.iter())
-        {
-            let cache_object = CacheChunkObject {
-                transform: object.transform,
-                id: object.id,
-                color: object.color,
-                pointer: StructurePointer::new(object.subdivided, false, 0),
-            };
-
-            let mut view = self.chunk_staging.write_buffer(
-                &mut encoder,
-                &self.chunk_objects,
-                *destination as u64 * size_of::<CacheChunkObject>() as u64,
-                NonZero::new(size_of::<CacheChunkObject>() as u64).unwrap(),
-            );
-
-            let temp_slice = [cache_object];
-            view.copy_from_slice(unsafe { as_raw_bytes(&temp_slice) });
-        }
-
-        self.update_indices();
-
-        self.chunk_staging.finish_and_recall_on_submit(&encoder);
-        self.queue.submit([encoder.finish()]);
-
-        destinations.iter().map(|d| ChunkHandle(*d)).collect()
-    }
-}

src/voxel_cache/cache.rs

@@ -13,10 +13,11 @@ use wgpu::ComputePipeline;
 use wgpu::Device;
 use wgpu::Queue;
 use wgpu::ShaderStages;
+use wgpu::util::DeviceExt;
 use wgpu::util::StagingBelt;
 
-use crate::voxel::gpu::StructurePointer;
-use crate::voxel::sparse::Color;
+use crate::voxel_cache::gpu::StructurePointer;
+use crate::voxel_cache::sparse::Color;
 
 #[derive(Clone, Copy, Zeroable, Pod, Debug)]
 #[repr(C)]
@@ -36,6 +37,9 @@ pub struct RequestBuffer<const N: usize>
     pub cache_size: usize,
     pub request_buffer: Buffer,
     pub request_count_buffer: Buffer,
+    pub indirect_count_storage: Buffer,
+    pub element_sort_count_buffer: Buffer,
+    pub indirect_count: Buffer,
     pub sort_buffer: Buffer,
     pub reset_pipeline: ComputePipeline,
     pub sort_pipeline: ComputePipeline,
@@ -74,6 +78,28 @@ where
             mapped_at_creation: false,
         });
 
+
+        let indirect_count = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("indirect_count_buffer"),
+            usage: BufferUsages::INDIRECT | BufferUsages::COPY_DST,
+            size: size_of::<wgpu::util::DispatchIndirectArgs>() as u64,
+            mapped_at_creation: false,
+        });
+
+        let indirect_count_storage = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("indirect_count_storage_buffer"),
+            usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
+            size: size_of::<wgpu::util::DispatchIndirectArgs>() as u64,
+            mapped_at_creation: false,
+        });
+
+        let element_sort_count_buffer = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("indirect_count_storage_buffer"),
+            usage: BufferUsages::STORAGE,
+            size: size_of::<u32>() as u64,
+            mapped_at_creation: false,
+        });
+
         // One element number per child entry
         // One number to identify node, one to identify sub child
         let sort_buffer_count = cache_size * N * N * N;
@@ -128,6 +154,26 @@ where
                     },
                     count: None,
                 },
+                wgpu::BindGroupLayoutEntry {
+                    binding: 3,
+                    visibility: ShaderStages::COMPUTE,
+                    ty: wgpu::BindingType::Buffer {
+                        ty: wgpu::BufferBindingType::Storage { read_only: false },
+                        has_dynamic_offset: false,
+                        min_binding_size: None,
+                    },
+                    count: None,
+                },
+                wgpu::BindGroupLayoutEntry {
+                    binding: 4,
+                    visibility: ShaderStages::COMPUTE,
+                    ty: wgpu::BindingType::Buffer {
+                        ty: wgpu::BufferBindingType::Storage { read_only: false },
+                        has_dynamic_offset: false,
+                        min_binding_size: None,
+                    },
+                    count: None,
+                },
             ],
         });
 
@@ -147,6 +193,14 @@ where
                     binding: 2,
                     resource: request_count_buffer.as_entire_binding(),
                 },
+                BindGroupEntry {
+                    binding: 3,
+                    resource: indirect_count_storage.as_entire_binding(),
+                },
+                BindGroupEntry {
+                    binding: 4,
+                    resource: element_sort_count_buffer.as_entire_binding(),
+                },
             ],
         });
 
@@ -181,6 +235,8 @@ where
                             @group(0) @binding(0) var<storage, read_write> request_buffer: array<RequestElement>;
                             @group(0) @binding(1) var<storage, read_write> sort_indirection: array<SortElement>; 
                             @group(0) @binding(2) var<storage, read_write> request_count: u32; 
+                            @group(0) @binding(3) var<storage, read_write> indirect_count: vec3<u32>;
+                            @group(0) @binding(4) var<storage, read_write> element_sort_count: u32; 
 
                             fn big_fusion(index: u32, phase_size: u32) -> vec2<u32>
                             {{
@@ -218,30 +274,24 @@ where
                                  let index = global_invocation_id.x;
                                  let total = arrayLength(&sort_indirection);
 
-                                 /*
-                                    let request_count_round_up =
-                                    ((request_count / {children_count}) + select(u32(0), u32(1), request_count % {children_count} != 0)) * {children_count};
-                                 if(index * 2 >= request_count_round_up)
-                                 {{ return; }}
-                                 */
-
                                  let sub_phase = (both_phase >> 16) & 0xFFFF;
                                  let phase = both_phase & 0xFFFF;
                                  let phase_total_width = u32(1 << (phase + 1));
 
-                                 request_count = 0;
-
 
                                  // Check if phase is last
                                  /*
-                                 let phase_count = u32(ceil(log2(f32(request_count))));
-                                 if(phase > phase_count)
+                                 if(phase == sub_phase)
                                  {{
-                                      return;
+                                      let next_phase_width = phase_total_width * 2;
+                                      if(next_phase_width >= element_sort_count && phase_total_width >= element_sort_count)
+                                      {{
+                                        // This was the final phase, stop
+                                        indirect_count = vec3<u32>(0);
+                                      }}
                                  }}
                                  */
 
-                                 /*
                                  if(phase == 0)
                                  {{
                                     let a = index * 2;
@@ -249,7 +299,6 @@ where
                                     sort_indirection[a].child = a % {children_count};
                                     sort_indirection[b].child = b % {children_count};
                                  }}
-                                 */
 
                                  var swap_indices = vec2<u32>(0, 0);
 
@@ -265,7 +314,7 @@ where
                                  // Do swap
 
                                  //if(swap_indices.y >= request_count_round_up)
-                                 if(swap_indices.y >= total)
+                                 if(swap_indices.y >= element_sort_count)
                                  {{
                                     // Suppose that swap_indices.y is -inf, dont swap
                                     return;
@@ -313,7 +362,9 @@ where
 
                             @group(0) @binding(0) var<storage, read_write> request_buffer: array<RequestElement>;
                             @group(0) @binding(1) var<storage, read_write> sort_indirection: array<SortElement>; 
-                            @group(0) @binding(2) var<storage, read_write> request_counts: u32; 
+                            @group(0) @binding(2) var<storage, read_write> request_counts: atomic<u32>; 
+                            @group(0) @binding(3) var<storage, read_write> indirect_count: u32; 
+                            @group(0) @binding(4) var<storage, read_write> node_sort_count: u32; 
 
                             @compute
                             @workgroup_size(64)
@@ -321,57 +372,55 @@ where
                                     @builtin(global_invocation_id) global_invocation_id: vec3<u32>
                                 )
                             {{
-                                 let index = global_invocation_id.x;
-                                 let sindex = index * {children_count};
-                                 let len = arrayLength(&request_buffer);
-                                 if(index >= len)
-                                 {{
+                                let index = global_invocation_id.x;
+                                let len = arrayLength(&request_buffer);
+                                if(index > len) {{ return; }}
+                                let sindex = index * {children_count};
+                                 
+                                if(phase == 0)
+                                {{
+                                    var count = 0;
+                                    for(var i = 0; i < {children_count}; i++)
+                                    {{
+                                        count += select(0, 1, request_buffer[index].children[i] != 0);
+                                    }}
+
+                                    sort_indirection[sindex].child = select(u32(0), u32(1), count != 0);
+                                    if count != 0
+                                    {{
+                                        atomicAdd(&request_counts, 1);
+                                    }}
                                     return;
-                                 }}
+                                }}
 
-                                 if(phase == u32(0xFFFFFFFF))
-                                 {{
-                                    request_counts = 0;
-                                    var children_sum = u32(0);
-                                    for(var i = u32(0); i < {children_count}; i++)
-                                    {{
-                                        children_sum += request_buffer[index].children[i];
-                                    }}
-                                    sort_indirection[sindex].child = select(u32(0), u32(1), children_sum != 0);
-                                    return;
-                                 }}
+                                if(phase == 0xFFFFFFFF)
+                                {{
+                                    // Double buffering bring back
+                                    sort_indirection[sindex].child = sort_indirection[sindex].node;
+                                }}
 
-                                 if(phase == u32(0x0FFFFFFF))
-                                 {{
-                                        sort_indirection[sindex].child = sort_indirection[sindex].node;
-                                        return;
-                                 }}
-
-                                 let phase_offset = u32(1) << phase;
-                                 let gather_offset = phase_offset * u32({children_count});
-                                 if(index >= phase_offset)
-                                 {{
-                                    // Double buffering
-                                    if(phase % 2 == 0)
-                                    {{
-                                        sort_indirection[sindex].node = sort_indirection[sindex].child + sort_indirection[sindex - gather_offset].child;
-                                    }}else
-                                    {{
-                                        sort_indirection[sindex].child = sort_indirection[sindex].node + sort_indirection[sindex - gather_offset].node;
-                                    }}
-                                 }}
-                                 else
-                                 {{ 
-                                    // Double buffering
-                                    if(phase % 2 == 0)
-                                    {{
-                                        sort_indirection[sindex].node = sort_indirection[sindex].child;
-                                    }}else
-                                    {{
-                                        sort_indirection[sindex].child = sort_indirection[sindex].node;
-                                    }}
-                                 }}
+                                // Phase is not zero, running sum part
+                                let running_sum_phase = phase - 1;
+                                let running_sum_offset = u32((1 << running_sum_phase) * {children_count});
 
+                                var add = u32(0);
+                                // Double buffering
+                                if(running_sum_phase % 2 == 0)
+                                {{
+                                     if(sindex >= running_sum_offset)
+                                     {{ 
+                                        add = sort_indirection[sindex - running_sum_offset].child;
+                                     }}
+                                     sort_indirection[sindex].node = sort_indirection[sindex].child + add;
+                                }}
+                                else
+                                {{
+                                     if(sindex >= running_sum_offset)
+                                     {{ 
+                                        add = sort_indirection[sindex - running_sum_offset].node;
+                                     }}
+                                     sort_indirection[sindex].child = sort_indirection[sindex].node + add;
+                                }}
                             }}
                         ")
                         .into(),
@@ -406,6 +455,8 @@ where
                             @group(0) @binding(0) var<storage, read_write> request_buffer: array<RequestElement>;
                             @group(0) @binding(1) var<storage, read_write> sort_indirection: array<SortElement>; 
                             @group(0) @binding(2) var<storage, read_write> request_counts: atomic<u32>; 
+                            @group(0) @binding(3) var<storage, read_write> indirect_count: vec3<u32>;
+                            @group(0) @binding(4) var<storage, read_write> element_sort_count: u32; 
 
                             @compute
                             @workgroup_size(64)
@@ -416,27 +467,28 @@ where
                                  let index = global_invocation_id.x;
                                  let sindex = index * {children_count};
                                  let len = arrayLength(&request_buffer);
-                                 if(index >= len)
+
+                                 element_sort_count = sort_indirection[(len - 1) * {children_count}].child * {children_count};
+                                 let invocation_count = (element_sort_count / 2) + select(u32(0), u32(1), element_sort_count % 2 != 0);
+                                 indirect_count = vec3(
+                                     (invocation_count / 64) + select(u32(0), u32(1), invocation_count % 64 != 0),
+                                     1, 1
+                                 );
+                                 let destination_index = sort_indirection[sindex].child - 1;
+
+                                 var count = u32(0);
+                                 for(var i = 0; i < {children_count}; i++)
                                  {{
-                                    return;
+                                      count += request_buffer[index].children[i];
                                  }}
 
-                                // Check if keep
-                                var node_request_count = u32(0);
-                                for(var i = 0; i < {children_count}; i++)
-                                {{
-                                    node_request_count += select(u32(0), u32(1), request_buffer[index].children[i] != 0);
-                                }}
-                                let keep = node_request_count != 0;
-                                atomicAdd(&request_counts, node_request_count);
-                                if(keep)
-                                {{
-                                    let running_sum_value = sort_indirection[sindex].child;
-                                    for(var i = u32(0); i < {children_count}; i++)
+                                 if(count != 0) // Keep ?
+                                 {{
+                                    for(var i = u32(0); i < u32({children_count}); i++)
                                     {{
-                                        sort_indirection[running_sum_value - 1 + i].node = index;
+                                        sort_indirection[destination_index * u32({children_count}) + i].node = index;
                                     }}
-                                }}
+                                 }}
                             }}
                         ")
                         .into(),
@@ -462,7 +514,8 @@ where
 
                             @group(0) @binding(0) var<storage, read_write> request_buffer: array<RequestElement>;
                             @group(0) @binding(1) var<storage, read_write> _ignore: array<u32>; 
-                            @group(0) @binding(2) var<storage, read_write> request_count: atomic<u32>; 
+                            @group(0) @binding(2) var<storage, read_write> request_count: u32;
+                            @group(0) @binding(4) var<storage, read_write> node_sort_count: u32; 
 
                             @compute
                             @workgroup_size(64)
@@ -479,12 +532,15 @@ where
                                     {{
                                         for(var i = 0; i < {children_count}; i += 1)
                                         {{
+                                             /*
                                             if(request_buffer[index].children[i] != 0)
                                             {{
                                                 atomicAdd(&request_count, 1);
                                             }}
+                                            */
                                             request_buffer[index].children[i] = 0;
                                         }}
+                                        request_count = 0;
                                     }}
                             }}
                         ")
@@ -500,6 +556,9 @@ where
             cache_size,
             request_buffer,
             request_count_buffer,
+            indirect_count,
+            indirect_count_storage,
+            element_sort_count_buffer,
             running_sum_pipeline,
             sort_buffer,
             sort_pipeline,
@@ -542,58 +601,71 @@ where
 
     pub fn sort_requests(&self, encoder: &mut CommandEncoder)
     {
-        let mut compute_pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
-            label: Some("request_buffer_sort_compute_pass"),
+        let mut compaction_compute_pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
+            label: Some("request_buffer_compaction_compute_pass"),
             timestamp_writes: None,
         });
 
-        compute_pass.set_bind_group(0, Some(&self.bindgroup), &[]);
+        compaction_compute_pass.set_bind_group(0, Some(&self.bindgroup), &[]);
         let request_element_count = self.cache_size; // Each child slot is sorted
         let workgroups_invocations = request_element_count.div_ceil(64);
 
-        /*
         // = Perform list compaction
         
         // == Running sum
-        compute_pass.set_pipeline(&self.running_sum_pipeline);
+        compaction_compute_pass.set_pipeline(&self.running_sum_pipeline);
 
         // Phase 0: Put ones in correct location
-        compute_pass.set_immediates(0, bytes_of(&u32::MAX));
-        compute_pass.dispatch_workgroups(workgroups_invocations as u32, 1, 1);
+        compaction_compute_pass.set_immediates(0, bytes_of(&0));
+        compaction_compute_pass.dispatch_workgroups(workgroups_invocations as u32, 1, 1);
         // Phase _: running sum
         // Running sum phase
-        for i in 0..(request_element_count.next_power_of_two().ilog2())
+        let running_sum_steps = request_element_count.next_power_of_two().ilog2();
+        for i in 1..=running_sum_steps
         {
-            compute_pass.set_immediates(0, bytes_of(&i));
-            compute_pass.dispatch_workgroups(workgroups_invocations as u32, 1, 1);
+            compaction_compute_pass.set_immediates(0, bytes_of(&i));
+            compaction_compute_pass.dispatch_workgroups(workgroups_invocations as u32, 1, 1);
         }
 
-        if !request_element_count.next_power_of_two().ilog2().is_multiple_of(2)
+        if !running_sum_steps.is_multiple_of(2)
         {
-            compute_pass.set_immediates(0, bytes_of(&0x0FFFFFFF));
-            compute_pass.dispatch_workgroups(workgroups_invocations as u32, 1, 1);
+            // Bring back double buffer
+            compaction_compute_pass.set_immediates(0, bytes_of(&0xFFFFFFFF_u32));
+            compaction_compute_pass.dispatch_workgroups(workgroups_invocations as u32, 1, 1);
         }
 
         // == Stream compaction
-        compute_pass.set_pipeline(&self.compaction_pipeline);
-        compute_pass.dispatch_workgroups(workgroups_invocations as u32, 1, 1);
-        */
+        compaction_compute_pass.set_pipeline(&self.compaction_pipeline);
+        compaction_compute_pass.dispatch_workgroups(workgroups_invocations as u32, 1, 1);
+        drop(compaction_compute_pass);
+
+        // == Copy count into indirect buffer
+
+
+        // = Sort
 
         // == Bitonic sort
-        compute_pass.set_pipeline(&self.sort_pipeline);
         let sort_element_count = self.cache_size * N * N * N;
         let phases_upper_bound = sort_element_count.next_power_of_two().ilog2();
 
         // Phase 0 dispatch all
         for i in 0..=phases_upper_bound
         {
+            encoder.copy_buffer_to_buffer(&self.indirect_count_storage, 0, &self.indirect_count, 0, Some(size_of::<wgpu::util::DispatchIndirectArgs>() as u64));
+            let mut sort_compute_pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
+                label: Some(format!("request_buffer_sort_compute_pass_{}", 0).as_str()),
+                timestamp_writes: None,
+            });
+            sort_compute_pass.set_bind_group(0, Some(&self.bindgroup), &[]);
+            sort_compute_pass.set_pipeline(&self.sort_pipeline);
             for j in 0..=i
             {
-                compute_pass.set_immediates(0, bytemuck::bytes_of(
+                sort_compute_pass.set_immediates(0, bytemuck::bytes_of(
                         &(i | (j << 16))
                         ));
-                compute_pass.dispatch_workgroups(sort_element_count.div_ceil(2).div_ceil(64) as u32, 1, 1);
+                sort_compute_pass.dispatch_workgroups_indirect(&self.indirect_count, 0);
             }
+            drop(sort_compute_pass);
         }
     }
 }
@@ -1631,7 +1703,7 @@ where
                                     }}
 
                                     let overwritten_element = lru_list[index];
-                                    if(parameters.write_pointers == 0)
+                                    if(parameters.write_pointers == 0 && usage_buffer[overwritten_element] != parameters.frame_timestamp)
                                     {{
                                         // Phase 1
                                         // Copy into cache page
@@ -1645,7 +1717,8 @@ where
                                         // Mark dirty/correct timestamp
                                         usage_buffer[overwritten_element] = parameters.frame_timestamp + 1;
                                     }}
-                                    else
+
+                                    if(parameters.write_pointers != 0 && usage_buffer[overwritten_element] != parameters.frame_timestamp)
                                     {{
                                         // Phase 2
                                         
@@ -1863,14 +1936,10 @@ where
         // Sort requests, reset request buffers to count requests
         self.usage_buffer.sort_usage(encoder);
         self.request_buffer.sort_requests(encoder);
-        self.request_buffer.reset_requests(encoder);
 
         self.structure_table
             .request_buffer_mut()
             .sort_requests(encoder);
-        self.structure_table
-            .request_buffer_mut()
-            .reset_requests(encoder);
 
         // Sorted requests are now in the group
         let user_target_bindgroup = self.device.create_bind_group(&wgpu::BindGroupDescriptor {
@@ -1899,6 +1968,12 @@ where
         let shader_invocation_count = request_target.size() / size_of::<CacheNodeRequest>() as u64;
         let workgroup_invocation_count = shader_invocation_count.div_ceil(64);
         write_requests_pass.dispatch_workgroups(workgroup_invocation_count as u32, 1, 1);
+        drop(write_requests_pass);
+
+        self.request_buffer.reset_requests(encoder);
+        self.structure_table
+            .request_buffer_mut()
+            .reset_requests(encoder);
     }
 
     pub fn current_timestamp(&self) -> u32

src/voxel_cache/gpu.rs

@@ -1,7 +1,7 @@
 use bytemuck::Pod;
 use bytemuck::Zeroable;
 
-use crate::voxel::sparse::Color;
+use crate::voxel_cache::sparse::Color;
 
 #[derive(Clone, Copy, Pod, Zeroable)]
 #[repr(transparent)]

src/voxel_cache/pipeline.rs

@@ -0,0 +1,68 @@
+use std::num::NonZero;
+
+use bytemuck::cast_slice;
+use crevice::std140::AsStd140;
+use glam::Mat4;
+use wgpu::BindGroup;
+use wgpu::BindGroupDescriptor;
+use wgpu::BindGroupEntry;
+use wgpu::BindGroupLayout;
+use wgpu::Buffer;
+use wgpu::BufferUsages;
+use wgpu::CommandEncoder;
+use wgpu::CommandEncoderDescriptor;
+use wgpu::ComputePass;
+use wgpu::ComputePassDescriptor;
+use wgpu::ComputePipeline;
+use wgpu::Device;
+use wgpu::Operations;
+use wgpu::Queue;
+use wgpu::RenderPipeline;
+use wgpu::ShaderModuleDescriptor;
+use wgpu::ShaderStages;
+use wgpu::TextureFormat;
+use wgpu::TextureView;
+use wgpu::VertexBufferLayout;
+use wgpu::util::StagingBelt;
+
+use crate::as_raw_bytes;
+use crate::camera::Camera;
+use crate::voxel_cache::cache::ColorPoolElement;
+use crate::voxel_cache::cache::LocationPoolElement;
+use crate::voxel_cache::cache::RequestBufferElement;
+use crate::voxel_cache::cache::StructurePoolElement;
+use crate::voxel_cache::gpu::StructurePointer;
+use crate::voxel_cache::sparse::Color;
+
+// Represents a chunk to be rendered by the voxel pipeline
+#[derive(Clone, Copy)]
+#[repr(C)]
+pub struct ChunkObject
+{
+    // Chunk object transform
+    pub transform: Mat4,
+
+    // Chunk data
+    pub color: Color,
+    pub subdivided: bool,
+
+    // Producer specific data
+    pub id: u32,
+}
+
+#[derive(Clone, Copy)]
+#[repr(C)]
+pub struct CacheChunkObject
+{
+    // Chunk object transform
+    transform: Mat4,
+
+    // Chunk data
+    color: Color,
+
+    // Producer specific data
+    id: u32,
+
+    // Pointer into cache
+    pointer: StructurePointer,
+}

src/voxel_cache/sparse.rs

@@ -5,8 +5,8 @@ use bytemuck::Pod;
 use bytemuck::Zeroable;
 use itertools::Itertools;
 
-use crate::voxel::gpu::ExplicitNTreeNode;
-use crate::voxel::gpu::StructurePointer;
+use crate::voxel_cache::gpu::ExplicitNTreeNode;
+use crate::voxel_cache::gpu::StructurePointer;
 
 #[derive(Debug, Clone, Copy, PartialEq, Pod, Zeroable)]
 #[repr(C)]