Initial

.gitignore

@@ -1,2 +1,9 @@
 /target
 Cargo.lock
+
+
+# Added by cargo
+#
+# already existing elements were commented out
+
+#/target

Cargo.toml

@@ -1,17 +1,18 @@
 [package]
-name = "wgpu-template"
+name = "vxls"
 version = "0.1.0"
 edition = "2024"
 
 [dependencies]
-anyhow = "1.0.100"
-bytemuck = "1.24.0"
-cgmath = "0.18.0"
-crevice = {version = "0.18.0", features = ["cgmath"]}
-egui = "0.33.3"
-egui-wgpu = "0.33.3"
-egui-winit = "0.33.3"
-env_logger = "0.11.8"
-pollster = "0.4.0"
-wgpu = "27.0.1"
-winit = "0.30.12"
+bytemuck = {version = "1.25.2", features = ["derive"]}
+crevice = {version = "0.20.1", features = ["glam"]}
+egui = "0.36.1"
+egui-wgpu = "0.36.1"
+egui-winit = "0.36.1"
+env_logger = "0.11.11"
+glam = "0.33.5"
+itertools = "0.15.0"
+pollster = "1.0.1"
+rand = "0.10.2"
+wgpu = "30"
+winit = "0.30.13"

shaders/cache.wgsl

@@ -0,0 +1,334 @@
+struct VoxelStructureNode
+{
+    pointers: array<u32, 64>
+}
+
+struct VoxelColorNode
+{
+    pointers: array<vec4<f32>, 64>
+}
+
+struct CacheNodeRequest
+{
+    location_info: u32,
+    request_count: u32
+}
+
+struct CacheNodeUsageBuffer
+{
+    touch_time: u32,
+    parent: u32,
+}
+
+// Rendering kernels
+struct ChunkInfo
+{
+    mvp: mat4x4<f32>,
+    eye_pos: vec3<f32>,
+    root_color: vec4<f32>,
+    root_subdiv: u32,
+    frame_timestamp: u32,
+}
+
+struct StructureNode
+{
+    pointers: array<u32, 64>
+}
+
+struct RequestElement
+{
+    pointers: array<atomic<u32>, 64>
+}
+
+struct ColorNode
+{
+    colors: array<vec4<f32>, 64>
+}
+
+struct VertexOutput
+{
+    @builtin(position) postion: vec4<f32>,
+    @location(0) world_loc: vec3<f32>,
+    @location(1) cam_pos: vec3<f32>,
+}
+
+var<immediate> constants: ChunkInfo;
+//var<push_constant> constants: ChunkInfo;
+
+@vertex
+fn chunk(@builtin(vertex_index) index: u32) -> VertexOutput
+{
+    let cube_vertices = array<vec3<f32>, 8>(
+        vec3<f32>(0., 0., 0.),
+        vec3<f32>(0., 0., 1.),
+        vec3<f32>(1., 0., 1.),
+        vec3<f32>(1., 0., 0.),
+
+        vec3<f32>(0., 1., 0.),
+        vec3<f32>(0., 1., 1.),
+        vec3<f32>(1., 1., 1.),
+        vec3<f32>(1., 1., 0.),
+    );
+
+    let cube_faces = array<u32, 24>(
+        // Bottom face
+        1, 0, 2, 3,
+
+        // Top face
+        4, 5, 7, 6,
+
+        // Side faces
+        0, 1, 4, 5,
+        1, 2, 5, 6,
+        2, 3, 6, 7,
+        3, 0, 7, 4,
+    );
+
+    let quad_index = index / (3 * 2);
+    let triangle_index = index % (3 * 2);
+    let triangle_map = array<u32, 6>(
+        0, 1, 2, 1, 3, 2
+    );
+
+    let vertex = cube_vertices[cube_faces[quad_index * 4 + triangle_map[triangle_index]]];
+    let output_vertex = constants.mvp * vec4<f32>(vertex, 1.0f);
+
+    var output: VertexOutput;
+    output.postion = output_vertex;
+    output.world_loc = vertex;
+    output.cam_pos = constants.eye_pos;
+
+    //let output = vec4<f32>(vertex, 1.0f);
+    return output;
+}
+
+@group(0) @binding(0) var<storage, read> structure_buffer: array<StructureNode>;
+@group(0) @binding(1) var<storage, read> color_buffer: array<ColorNode>;
+@group(0) @binding(2) var<storage, read_write> request_buffer: array<RequestElement>;
+@group(0) @binding(3) var<storage, read_write> usage_buffer: array<u32>;
+
+
+fn box_inter(pos: vec3<f32>, ray_dir: vec3<f32>, box_min: vec3<f32>, box_max: vec3<f32>) -> vec2<f32>
+{
+    let box_min_t = (box_min - pos) / ray_dir;
+    let box_max_t = (box_max - pos) / ray_dir;
+
+    let near_ts = min(box_min_t, box_max_t);
+    let far_ts = max(box_min_t, box_max_t);
+
+    let far_t = min(min(far_ts.x, far_ts.y), far_ts.z);
+    let near_t = max(max(near_ts.x, near_ts.y), near_ts.z);
+
+    return vec2(near_t, far_t);
+}
+
+fn sdf(voxel: vec3<i32>) -> bool
+{
+    let len = length(vec3<f32>(voxel) - vec3(128)) / 128.;
+    return len <= 1.;
+}
+
+fn min_vec(x: vec3<f32>) -> f32
+{
+    return min(x.x, min(x.y, x.z));
+}
+
+fn min_mask(x: vec3<f32>) -> vec3<bool>
+{
+    let min = min(x.x, min(x.y, x.z));
+
+    return vec3<bool>(min == x.x, min == x.y, min == x.z);
+}
+
+fn node_subdivided(node: u32) -> bool
+{
+    return (node >> 31) != 0;
+}
+
+fn node_pointer_valid(node: u32) -> bool
+{
+    return ((node >> 30) & 1) != 0;
+}
+
+fn node_pointer(node: u32) -> u32
+{
+    return node & 0x3FFFFFFF;
+}
+
+fn voxel_from_wall(position: vec3<f32>, ray_dir: vec3<f32>) -> vec3<i32>
+{
+    let integers = round(position);
+    let wall_mask = min_mask(abs(position - vec3<f32>(integers)));
+    let offsets = select(vec3<f32>(-0.5), vec3<f32>(0.5), ray_dir > vec3(0.));
+    return vec3<i32>(floor(position + select(vec3<f32>(0.), offsets, wall_mask)));
+}
+
+fn traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_color: vec4<f32>, root_subdiv: bool) -> vec4<f32>
+{
+    // Simple FVT
+    let t_off = abs(1. / ray_dir);
+
+    // Start location
+    let voxel_dir = select(vec3(-1), vec3(1), ray_dir >= vec3(0.));
+    var pos_origin = clamp(ray_origin * 256., vec3(0.), vec3(256. - 1.));
+    var voxel = vec3<i32>(pos_origin);
+    var last_voxel = voxel;
+
+    let wall_offset = select(vec3(0), vec3(1), ray_dir > vec3(0.));
+
+    var dfs_stack = array<u32, 5>(0, 0, 0, 0, 0);
+
+        // Current depth of the node we are exploring
+        var current_depth = 0;
+
+        // Index of the current node's data
+        var current_node = u32(0);
+
+        // Current node size
+        var node_size = 4 * 4 * 4 * 4; // 128
+
+        // Size of a child of this node
+        var child_size = node_size / 4;
+
+        // Lut of the node_size per depth
+        var node_size_lut = array<i32, 5>(
+                4 * 4 * 4 * 4,
+                4 * 4 * 4,
+                4 * 4,
+                4,
+                1,
+                );
+
+    let depth_limit = 1;
+    for(var iter = 0; iter < 256; iter ++)
+    {
+
+        // Our ray is currently touching a voxel.
+        // Descend to the lowest node that contains this voxel
+
+        // Position of the child we are in
+        var child_pos = (vec3<u32>(voxel) >> vec3<u32>((4 - u32(current_depth + 1)) * 2)) & vec3<u32>(3); // Hardcode for 4-tree
+        var child_index = child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4;
+
+        // Current node has been used. report
+        usage_buffer[current_node] = constants.frame_timestamp;
+        
+        while(
+            node_subdivided(structure_buffer[current_node].pointers[child_index]) 
+            && current_depth < depth_limit)
+        {
+            if(!node_pointer_valid(structure_buffer[current_node].pointers[child_index]))
+            {
+                atomicAdd(&request_buffer[current_node].pointers[child_index], 1);
+                break;
+            }
+            // Child node is subdivided, we go in, save position in stack
+            current_node = node_pointer(structure_buffer[current_node].pointers[child_index]);
+            usage_buffer[current_node] = constants.frame_timestamp;
+            current_depth += 1;
+            dfs_stack[current_depth] = current_node;
+            node_size = node_size_lut[current_depth];
+
+            child_pos = (vec3<u32>(voxel) >> vec3<u32>((4 - u32(current_depth + 1)) * 2)) & vec3<u32>(3); // Hardcode for 4-tree
+            child_index = child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4;
+            child_size = node_size / 4;
+        }
+
+
+        // At this point current_depth is the depth of the node that contains the voxel
+        // child_pos and child_index relate to the specific child of the node that contains this voxel
+        // It is guaranteed that the child is leave
+
+        // Check current leave's color
+        let color = color_buffer[current_node].colors[child_index];
+        if(color.w != 0.) // Not transparent
+        {
+            let k = child_pos.x + child_pos.y + child_pos.z;
+            let x = select(0.5, 1., k % 2 == 0);
+            return x * color / f32(current_depth);
+        }
+
+        // Voxel and whole child containing it is empty
+        
+        // Perform a step through the children of the node
+        let child_position = (voxel / child_size) * child_size;
+        let far_corner = child_position + wall_offset * child_size;
+        let far_ts = (vec3<f32>(far_corner) - pos_origin) / ray_dir; // TODO: Turn into fma
+        let far_t = min(min(far_ts.x, far_ts.y), far_ts.z);
+
+        let next_child_min = select(child_position, child_position + wall_offset * child_size, vec3(far_t) == far_ts);
+        let next_child_max = next_child_min + vec3(child_size);
+        
+        // The ray (far_t) is now touching the new child to explore
+        // Find out which actual voxel we are touching
+        let previous_voxel = voxel;
+        let float_voxel = clamp(pos_origin + far_t * ray_dir, vec3<f32>(next_child_min), vec3<f32>(next_child_max));
+        /*
+        voxel = vec3<i32>(
+            floor(
+            select(
+                float_voxel - vec3(0.5),
+                float_voxel + vec3(0.5),
+                ray_dir > vec3(0.)
+            ))
+        );
+        */
+        //voxel = vec3<i32>(round(float_voxel));
+        voxel = voxel_from_wall(float_voxel, ray_dir);
+        if(any(voxel < vec3(0)) || any(voxel >= vec3(256)))
+        {
+            //return vec4(f32(iter) / 100.);
+            discard;
+        }
+
+        // We touched a voxel as if we explored blocks sized by the child size of the current node.
+        // But we might have exited the current node.
+        
+        // If this is the case we have to walk back up the tree
+        // And then back down to the next node over
+
+        // As such we find the lowest ancestor that can contain both the privous voxel (in node) and the new voxel (out of node)
+        let bit_diffs = voxel ^ previous_voxel;
+        let bit_diffs_lowest = bit_diffs.x | bit_diffs.y | bit_diffs.z;
+
+        let flb = ((countLeadingZeros(bit_diffs_lowest) - 24) / 2);
+        let common_depth = flb;
+
+        current_depth = common_depth;
+        node_size = node_size_lut[current_depth];
+        child_size = node_size / 4;
+        current_node = dfs_stack[current_depth];
+        
+        // Figure out current voxel position
+        //voxel = vec3<i32>(ray_origin + ray_dir * t); 
+
+    }
+    return vec4<f32>(1., 0., 1., 1.);
+
+}
+
+@fragment
+fn fragment(in: VertexOutput) -> @location(0) vec4<f32>
+{
+    let ray_dir = normalize(in.world_loc - in.cam_pos);
+    let interp = box_inter(in.cam_pos, ray_dir, vec3(0.), vec3(1));
+    let ray_origin = in.cam_pos + ray_dir * (max(0., interp.x));
+
+    return traverse(ray_dir, ray_origin, constants.root_color, constants.root_subdiv != 0);
+    return vec4(interp.y / 10.);
+}
+
+// Cache managment kernels
+
+// Request buffer managment
+// Compaction
+@workgroup_size(64)
+fn main(
+    @builtin(global_invocation_id) gid: vec3<u32>
+)
+{
+    let cache_size = 1024;
+    let index = gid.x;
+
+    for()
+}

shaders/example.slang

@@ -0,0 +1,2 @@
+module example;
+

shaders/voxel.wgsl

@@ -0,0 +1,390 @@
+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>,
+}
+
+struct ChunkImmediate
+{
+    view_proj: mat4x4<f32>,
+    cam_pos: vec3<f32>,
+    frame_timestamp: u32
+}
+
+var<immediate> constants: ChunkImmediate;
+//var<push_constant> constants: ChunkInfo;
+
+struct CacheChunkObject
+{
+    transform: mat4x4<f32>,
+    color: vec4<f32>,
+    id: u32,
+    pointer: u32
+}
+
+
+struct StructurePoolElement
+{
+    pointers: array<u32, 64>
+}
+
+struct RequestBufferElement
+{
+    requests: array<atomic<u32>, 64>
+}
+
+struct ColorPoolElement
+{
+    colors: array<vec4<f32>, 64>
+}
+
+struct LocationPoolElement
+{
+    structure_id: u32,
+    structure_locator: u32
+}
+
+struct SortedRequestsElement
+{
+    node: u32,
+    child: u32
+}
+
+@group(0) @binding(0) var<storage, read_write> structure_pool: array<StructurePoolElement>;
+@group(0) @binding(1) var<storage, read_write> color_pool: array<ColorPoolElement>;
+@group(0) @binding(2) var<storage, read_write> location_pool: array<LocationPoolElement>;
+@group(0) @binding(3) var<storage, read_write> request_buffer: array<RequestBufferElement>;
+@group(0) @binding(4) var<storage, read_write> usage_buffer: array<atomic<u32>>;
+@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>>;
+
+@vertex
+fn chunk(@builtin(vertex_index) index: u32) -> VertexOutput
+{
+    let cube_vertices = array<vec3<f32>, 8>(
+        vec3<f32>(0., 0., 0.),
+        vec3<f32>(0., 0., 1.),
+        vec3<f32>(1., 0., 1.),
+        vec3<f32>(1., 0., 0.),
+
+        vec3<f32>(0., 1., 0.),
+        vec3<f32>(0., 1., 1.),
+        vec3<f32>(1., 1., 1.),
+        vec3<f32>(1., 1., 0.),
+    );
+
+    let cube_faces = array<u32, 24>(
+        // Bottom face
+        1, 0, 2, 3,
+
+        // Top face
+        4, 5, 7, 6,
+
+        // Side faces
+        0, 1, 4, 5,
+        1, 2, 5, 6,
+        2, 3, 6, 7,
+        3, 0, 7, 4,
+    );
+
+    let quad_index = index / (3 * 2);
+    let triangle_index = index % (3 * 2);
+    let triangle_map = array<u32, 6>(
+        0, 1, 2, 1, 3, 2
+    );
+
+
+    let vertex = cube_vertices[cube_faces[quad_index * 4 + triangle_map[triangle_index]]];
+    let output_vertex = constants.view_proj * vec4<f32>(vertex, 1.0f);
+
+    var output: VertexOutput;
+    output.postion = output_vertex;
+    output.color = vec4(1.);
+    output.chunk_index = 0;
+    output.cam_pos = constants.cam_pos;
+    output.world_pos = vertex;
+
+    //let output = vec4<f32>(vertex, 1.0f);
+    return output;
+}
+
+
+struct StructureElement
+{
+    children: array<u32, 64> 
+}
+
+struct ColorElement 
+{
+    children: array<vec4<f32>, 64> 
+}
+
+struct LocationElement 
+{
+    children: array<vec4<f32>, 64> 
+}
+
+struct RequestElement
+{
+    children: array<atomic<u32>, 64> 
+}
+
+fn box_inter(pos: vec3<f32>, ray_dir: vec3<f32>, box_min: vec3<f32>, box_max: vec3<f32>) -> vec2<f32>
+{
+    let box_min_t = (box_min - pos) / ray_dir;
+    let box_max_t = (box_max - pos) / ray_dir;
+
+    let near_ts = min(box_min_t, box_max_t);
+    let far_ts = max(box_min_t, box_max_t);
+
+    let far_t = min(min(far_ts.x, far_ts.y), far_ts.z);
+    let near_t = max(max(near_ts.x, near_ts.y), near_ts.z);
+
+    return vec2(near_t, far_t);
+}
+
+fn sdf(voxel: vec3<i32>) -> bool
+{
+    let len = length(vec3<f32>(voxel) - vec3(128)) / 128.;
+    return len <= 1.;
+}
+
+fn min_vec(x: vec3<f32>) -> f32
+{
+    return min(x.x, min(x.y, x.z));
+}
+
+fn min_mask(x: vec3<f32>) -> vec3<bool>
+{
+    let min = min(x.x, min(x.y, x.z));
+
+    return vec3<bool>(min == x.x, min == x.y, min == x.z);
+}
+
+fn node_subdivided(node: u32) -> bool
+{
+    return (node >> 31) != 0;
+}
+
+fn node_pointer_valid(node: u32) -> bool
+{
+    return ((node >> 30) & 1) != 0;
+}
+
+fn node_pointer(node: u32) -> u32
+{
+    return node & 0x3FFFFFFF;
+}
+
+fn voxel_from_wall(position: vec3<f32>, ray_dir: vec3<f32>) -> vec3<i32>
+{
+    let integers = round(position);
+    let wall_mask = min_mask(abs(position - vec3<f32>(integers)));
+    let offsets = select(vec3<f32>(-0.5), vec3<f32>(0.5), ray_dir > vec3(0.));
+    return vec3<i32>(floor(position + select(vec3<f32>(0.), offsets, wall_mask)));
+}
+
+fn traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_color: vec4<f32>, root_subdiv: bool) -> vec4<f32>
+{
+    // Simple FVT
+    let t_off = abs(1. / ray_dir);
+
+    // Start location
+    let voxel_dir = select(vec3(-1), vec3(1), ray_dir >= vec3(0.));
+    var pos_origin = clamp(ray_origin * 256., vec3(0.), vec3(256. - 1.));
+    var voxel = vec3<i32>(pos_origin);
+    var last_voxel = voxel;
+
+    let wall_offset = select(vec3(0), vec3(1), ray_dir > vec3(0.));
+
+    var dfs_stack = array<u32, 5>(0, 0, 0, 0, 0);
+
+        // Current depth of the node we are exploring
+        var current_depth = 0;
+
+        // Index of the current node's data
+        var current_node = u32(structure_table_pointer[0] & 0x3FFFFFFF);
+
+        // Current node size
+        var node_size = 4 * 4 * 4 * 4; // 128
+
+        // Size of a child of this node
+        var child_size = node_size / 4;
+
+        // Lut of the node_size per depth
+        var node_size_lut = array<i32, 5>(
+                4 * 4 * 4 * 4,
+                4 * 4 * 4,
+                4 * 4,
+                4,
+                1,
+                );
+
+    let depth_limit = 3;
+    for(var iter = 0; iter < 256; iter ++)
+    {
+
+        // Our ray is currently touching a voxel.
+        // Descend to the lowest node that contains this voxel
+
+        // Position of the child we are in
+        var child_pos = (vec3<u32>(voxel) >> vec3<u32>((4 - u32(current_depth + 1)) * 2)) & vec3<u32>(3); // Hardcode for 4-tree
+        var child_index = child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4;
+
+        // Current node has been used. report
+        usage_buffer[current_node] = constants.frame_timestamp;
+        
+        while(
+            node_subdivided(structure_pool[current_node].pointers[child_index]) 
+            && current_depth < depth_limit)
+        {
+            if(!node_pointer_valid(structure_pool[current_node].pointers[child_index]))
+            {
+                atomicAdd(&request_buffer[current_node].requests[child_index], 1);
+                break;
+            }
+            // Child node is subdivided, we go in, save position in stack
+            current_node = node_pointer(structure_pool[current_node].pointers[child_index]);
+            usage_buffer[current_node] = constants.frame_timestamp;
+            current_depth += 1;
+            dfs_stack[current_depth] = current_node;
+            node_size = node_size_lut[current_depth];
+
+            child_pos = (vec3<u32>(voxel) >> vec3<u32>((4 - u32(current_depth + 1)) * 2)) & vec3<u32>(3); // Hardcode for 4-tree
+            child_index = child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4;
+            child_size = node_size / 4;
+        }
+        usage_buffer[current_node] = constants.frame_timestamp;
+
+
+
+
+        // At this point current_depth is the depth of the node that contains the voxel
+        // child_pos and child_index relate to the specific child of the node that contains this voxel
+        // It is guaranteed that the child is leave
+
+        // Check current leave's color
+        let color = color_pool[current_node].colors[child_index];
+        if(color.w != 0.) // Not transparent
+        {
+            let k = child_pos.x + child_pos.y + child_pos.z;
+            let w = voxel.x + voxel.y + voxel.z;
+            let x = select(0.5, 1., k % 2 == 0) * select(0.8, 1., w % 2 == 0);
+
+            var div = 1;
+            var overlay = 1.;
+            for(var i = 1; i <= 4; i++)
+            {
+                let x = (voxel.x / div + voxel.y / div + voxel.z / div) % 2 == 0;
+                overlay -= select(0., 1. / (f32(i) * 2.5), x);
+                div *= 4;
+            }
+
+            return overlay * color;
+        }
+
+        // Voxel and whole child containing it is empty
+        
+        // Perform a step through the children of the node
+        let child_position = (voxel / child_size) * child_size;
+        let far_corner = child_position + wall_offset * child_size;
+        let far_ts = (vec3<f32>(far_corner) - pos_origin) / ray_dir; // TODO: Turn into fma
+        let far_t = min(min(far_ts.x, far_ts.y), far_ts.z);
+
+        let next_child_min = select(child_position, child_position + wall_offset * child_size, vec3(far_t) == far_ts);
+        let next_child_max = next_child_min + vec3(child_size);
+        
+        // The ray (far_t) is now touching the new child to explore
+        // Find out which actual voxel we are touching
+        let previous_voxel = voxel;
+        let float_voxel = clamp(pos_origin + far_t * ray_dir, vec3<f32>(next_child_min), vec3<f32>(next_child_max));
+        /*
+        voxel = vec3<i32>(
+            floor(
+            select(
+                float_voxel - vec3(0.5),
+                float_voxel + vec3(0.5),
+                ray_dir > vec3(0.)
+            ))
+        );
+        */
+        //voxel = vec3<i32>(round(float_voxel));
+        voxel = voxel_from_wall(float_voxel, ray_dir);
+        if(any(voxel < vec3(0)) || any(voxel >= vec3(256)))
+        {
+            //return vec4(f32(iter) / 100.);
+            discard;
+        }
+
+        // We touched a voxel as if we explored blocks sized by the child size of the current node.
+        // But we might have exited the current node.
+        
+        // If this is the case we have to walk back up the tree
+        // And then back down to the next node over
+
+        // As such we find the lowest ancestor that can contain both the privous voxel (in node) and the new voxel (out of node)
+        let bit_diffs = voxel ^ previous_voxel;
+        let bit_diffs_lowest = bit_diffs.x | bit_diffs.y | bit_diffs.z;
+
+        let flb = ((countLeadingZeros(bit_diffs_lowest) - 24) / 2);
+        let common_depth = flb;
+
+        current_depth = common_depth;
+        node_size = node_size_lut[current_depth];
+        child_size = node_size / 4;
+        current_node = dfs_stack[current_depth];
+        
+        // Figure out current voxel position
+        //voxel = vec3<i32>(ray_origin + ray_dir * t); 
+
+    }
+    return vec4<f32>(1., 0., 1., 1.);
+
+}
+
+@fragment
+fn fragment(in: VertexOutput) -> @location(0) vec4<f32>
+{
+    let ray_dir = normalize(in.world_pos - in.cam_pos);
+    let interp = box_inter(in.cam_pos, ray_dir, vec3(0.), vec3(1));
+    let ray_origin = in.cam_pos + ray_dir * (max(0., interp.x));
+
+    if(length(ray_origin) < 0.05)
+    {
+        return vec4(1., 0., 0., 1.);
+    }
+
+    let root_subdiv = ((structure_table_pointer[0] >> 31) & 1) != 0;
+    let pointer_valid = ((structure_table_pointer[0] >> 30) & 1) != 0;
+    if(!pointer_valid && root_subdiv)
+    {
+        atomicAdd(&structure_table_request_buffer[0], 1);
+    }
+    if(!pointer_valid)
+    {
+        return vec4(0., 1., 0., 1.);
+    }
+    return traverse(ray_dir, ray_origin, vec4(1., 1., 1., 1.), root_subdiv);
+    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.);
+}
+*/
+

shaders/wesl.toml

@@ -0,0 +1 @@
+root = "."

src/camera.rs

@@ -0,0 +1,120 @@
+use std::collections::HashSet;
+use std::f32::consts::PI;
+
+use glam::Mat4;
+use glam::Vec2;
+use glam::Vec3;
+use winit::event::WindowEvent;
+use winit::keyboard::KeyCode;
+
+pub struct Camera
+{
+    pub position: Vec3,
+    pitch: f32,
+    yaw: f32,
+
+    zfar: f32,
+    znear: f32,
+    fov: f32,
+    pub aspect: f32,
+
+    speed: f32,
+    pub pressed_keyset: HashSet<KeyCode>,
+}
+
+impl Default for Camera
+{
+    fn default() -> Self
+    {
+        Self {
+            position: Vec3::new(0., 0., 0.),
+            pitch: Default::default(),
+            yaw: Default::default(),
+            speed: 0.1,
+            pressed_keyset: Default::default(),
+
+            znear: 0.01,
+            zfar: 100.,
+            fov: PI * 100.0f32 / 180.0f32,
+            aspect: 1.,
+        }
+    }
+}
+
+impl Camera
+{
+    pub fn handle_input(&mut self, event: &WindowEvent)
+    {
+        if let WindowEvent::KeyboardInput { event, .. } = event
+        {
+            match (event.state, event.physical_key)
+            {
+                (winit::event::ElementState::Pressed, winit::keyboard::PhysicalKey::Code(c)) =>
+                {
+                    self.pressed_keyset.insert(c);
+                }
+                (winit::event::ElementState::Released, winit::keyboard::PhysicalKey::Code(c)) =>
+                {
+                    self.pressed_keyset.remove(&c);
+                }
+                _ =>
+                {}
+            }
+        }
+    }
+
+    pub fn update(&mut self)
+    {
+        let mut movement = Vec3::new(0., 0., 0.);
+        if self.pressed_keyset.contains(&KeyCode::KeyW)
+        {
+            movement.z += self.speed;
+        }
+        if self.pressed_keyset.contains(&KeyCode::KeyS)
+        {
+            movement.z -= self.speed;
+        }
+
+        // Left right
+        if self.pressed_keyset.contains(&KeyCode::KeyA)
+        {
+            movement.x += self.speed;
+        }
+        if self.pressed_keyset.contains(&KeyCode::KeyD)
+        {
+            movement.x -= self.speed;
+        }
+
+        let rot_movement = glam::Mat3::from_rotation_y(-self.yaw)
+            * glam::Mat3::from_rotation_x(-self.pitch)
+            * movement;
+        self.position -= rot_movement;
+    }
+
+    pub fn cursor_moved(&mut self, x: f32, y: f32)
+    {
+        const SENSITIVITY: f32 = 0.0004;
+        let position = Vec2::new(x, y);
+        let offset = position * SENSITIVITY;
+
+        self.yaw += offset.x;
+        self.pitch += offset.y;
+    }
+
+    pub fn view_proj(&self) -> Mat4
+    {
+        let view = glam::Mat4::from_translation(self.position)
+            * glam::Mat4::from_rotation_y(-self.yaw)
+            * glam::Mat4::from_rotation_x(-self.pitch);
+        let proj = glam::Mat4::perspective_rh(self.fov, self.aspect, self.znear, self.zfar);
+        OPENGL_TO_WGPU_MATRIX * proj * view.inverse()
+    }
+}
+
+#[rustfmt::skip]
+pub const OPENGL_TO_WGPU_MATRIX: glam::Mat4 = glam::Mat4::from_cols(
+    glam::Vec4::new(1.0, 0.0, 0.0, 0.0),
+    glam::Vec4::new(0.0, 1.0, 0.0, 0.0),
+    glam::Vec4::new(0.0, 0.0, 0.5, 0.0),
+    glam::Vec4::new(0.0, 0.0, 0.5, 1.0),
+);

src/egui_renderer.rs

@@ -1,58 +1,68 @@
-use egui::{Context, PaintCallbackInfo};
-use egui_wgpu::{CallbackResources, CallbackTrait, Renderer, RendererOptions, ScreenDescriptor};
+use egui::Context;
+use egui_wgpu::Renderer;
+use egui_wgpu::RendererOptions;
+use egui_wgpu::ScreenDescriptor;
 use egui_winit::State;
-use wgpu::{CommandEncoder, Device, Queue, RenderPass, TextureFormat, TextureView};
-use winit::{event::WindowEvent, window::Window};
+use wgpu::CommandEncoder;
+use wgpu::Device;
+use wgpu::Queue;
+use wgpu::RenderPassDescriptor;
+use wgpu::TextureFormat;
+use wgpu::TextureView;
+use winit::event::WindowEvent;
+use winit::window::Window;
 
-pub struct EguiState {
-    pub state: State,
-    pub renderer: Renderer,
-    pub frame_started: bool,
-    //pub msaa_texture: TextureView,
-    pub color_format: TextureFormat,
+pub struct EguiRenderer
+{
+    state: State,
+    renderer: Renderer,
+    frame_started: bool,
 }
 
-impl EguiState {
-    pub fn context(&self) -> &Context {
+impl EguiRenderer
+{
+    pub fn context(&self) -> &Context
+    {
         self.state.egui_ctx()
     }
 
-    pub fn new(device: &Device, output_color_format: TextureFormat, window: &Window) -> EguiState {
+    pub fn new(device: &Device, output_color_format: TextureFormat, window: &Window)
+    -> EguiRenderer
+    {
         let egui_context = Context::default();
 
-        let egui_state = egui_winit::State::new(
+        let egui_state = State::new(
             egui_context,
-            egui::ViewportId::ROOT,
+            egui::viewport::ViewportId::ROOT,
             &window,
             Some(window.scale_factor() as f32),
             None,
             Some(2 * 1024),
         );
 
-        let options = RendererOptions {
-            //msaa_samples: SAMPLE_COUNT,
-            //depth_stencil_format: Some(TextureFormat::Depth24PlusStencil8),
-            ..Default::default()
-        };
-        let egui_renderer = Renderer::new(device, output_color_format, options);
+        let egui_renderer = Renderer::new(device, output_color_format, RendererOptions::default());
 
-        EguiState {
+        EguiRenderer {
             state: egui_state,
             renderer: egui_renderer,
             frame_started: false,
-            color_format: output_color_format,
         }
     }
 
-    pub fn resize(&mut self, _device: &Device, _width: u32, _height: u32) {}
-
-    pub fn handle_event(&mut self, window: &Window, event: &WindowEvent) {
+    pub fn handle_input(&mut self, window: &Window, event: &WindowEvent)
+    {
         let _ = self.state.on_window_event(window, event);
     }
 
-    pub fn begin_frame(&mut self, window: &Window) {
-        let input = self.state.take_egui_input(window);
-        self.state.egui_ctx().begin_pass(input);
+    pub fn ppp(&mut self, v: f32)
+    {
+        self.context().set_pixels_per_point(v);
+    }
+
+    pub fn begin_frame(&mut self, window: &Window)
+    {
+        let raw_input = self.state.take_egui_input(window);
+        self.state.egui_ctx().begin_pass(raw_input);
         self.frame_started = true;
     }
 
@@ -64,14 +74,15 @@ impl EguiState {
         window: &Window,
         window_surface_view: &TextureView,
         screen_descriptor: ScreenDescriptor,
-    ) {
-        if !self.frame_started {
-            panic!("begin_frame must be called before end_frame_and_draw can be called!");
+    )
+    {
+        if !self.frame_started
+        {
+            panic!("begin_frame must be called before.");
         }
 
-        //self.ppp(screen_descriptor.pixels_per_point);
-
-        let full_output = self.state.egui_ctx().end_pass();
+        self.ppp(screen_descriptor.pixels_per_point);
+        let mut full_output = self.state.egui_ctx().end_pass();
 
         self.state
             .handle_platform_output(window, full_output.platform_output);
@@ -80,58 +91,41 @@ impl EguiState {
             .state
             .egui_ctx()
             .tessellate(full_output.shapes, self.state.egui_ctx().pixels_per_point());
-        for (id, image_delta) in &full_output.textures_delta.set {
+
+        for (id, image_delta) in &full_output.textures_delta.set
+        {
             self.renderer
-                .update_texture(device, queue, *id, image_delta);
+                .update_texture(device, queue, *id, &image_delta[0]);
         }
+
+        full_output.textures_delta.clear();
         self.renderer
             .update_buffers(device, queue, encoder, &tris, &screen_descriptor);
-        let rpass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
+
+        let rpass = encoder.begin_render_pass(&RenderPassDescriptor {
             color_attachments: &[Some(wgpu::RenderPassColorAttachment {
                 view: window_surface_view,
+                depth_slice: None,
                 resolve_target: None,
-                ops: egui_wgpu::wgpu::Operations {
-                    load: egui_wgpu::wgpu::LoadOp::Load,
+                ops: wgpu::Operations {
+                    load: wgpu::LoadOp::Load,
                     store: wgpu::StoreOp::Store,
                 },
-                depth_slice: None,
             })],
+            multiview_mask: None,
             depth_stencil_attachment: None,
-            label: Some("egui main render pass"),
             timestamp_writes: None,
+            label: Some("egui main render pass"),
             occlusion_query_set: None,
         });
 
         self.renderer
             .render(&mut rpass.forget_lifetime(), &tris, &screen_descriptor);
-        for x in &full_output.textures_delta.free {
-            self.renderer.free_texture(x)
+        for x in &full_output.textures_delta.free
+        {
+            self.renderer.free_texture(x);
         }
 
         self.frame_started = false;
     }
 }
-
-pub struct CallbackFn<P>
-where
-    P: Fn(PaintCallbackInfo, &mut RenderPass<'static>, &CallbackResources),
-{
-    //pub : FnOnce(&Device, &Queue, &ScreenDescriptor, &mut CommandEncoder, &mut CallbackResources) -> Vec<>
-    pub paint_fn: P,
-}
-
-impl<P> CallbackTrait for CallbackFn<P>
-where
-    P: Fn(PaintCallbackInfo, &mut RenderPass<'static>, &CallbackResources)
-        + std::marker::Sync
-        + std::marker::Send,
-{
-    fn paint(
-        &self,
-        info: egui::PaintCallbackInfo,
-        render_pass: &mut wgpu::RenderPass<'static>,
-        callback_resources: &egui_wgpu::CallbackResources,
-    ) {
-        (self.paint_fn)(info, render_pass, callback_resources)
-    }
-}

src/lib.rs

@@ -1,74 +0,0 @@
-pub mod egui_renderer;
-pub mod state;
-
-use std::sync::Arc;
-
-use winit::{
-    application::ApplicationHandler,
-    event::WindowEvent,
-    event_loop::{self, EventLoop},
-    window::Window,
-};
-
-use crate::state::State;
-
-pub fn run() -> anyhow::Result<()> {
-    env_logger::init();
-
-    let event_loop = EventLoop::with_user_event().build()?;
-    let mut app = App::default();
-    event_loop.run_app(&mut app)?;
-
-    Ok(())
-}
-
-// App struct
-#[derive(Default)]
-pub struct App {
-    state: Option<State>,
-}
-
-impl ApplicationHandler for App {
-    fn resumed(&mut self, event_loop: &event_loop::ActiveEventLoop) {
-        // Create window
-        let window = Arc::new(
-            event_loop
-                .create_window(
-                    Window::default_attributes()
-                        .with_title("Wgpu Template")
-                        .with_resizable(true),
-                )
-                .unwrap(),
-        );
-
-        let state = pollster::block_on(State::new(window.clone()));
-        self.state = Some(state);
-
-        window.request_redraw();
-    }
-
-    fn window_event(
-        &mut self,
-        event_loop: &event_loop::ActiveEventLoop,
-        _window_id: winit::window::WindowId,
-        event: winit::event::WindowEvent,
-    ) {
-        let state = self.state.as_mut().unwrap();
-        state.handle_event(&event);
-
-        match event {
-            WindowEvent::CloseRequested => {
-                event_loop.exit();
-            }
-            WindowEvent::RedrawRequested => {
-                state.render();
-                state.get_window().request_redraw();
-            }
-
-            WindowEvent::Resized(size) => {
-                state.resize(size);
-            }
-            _ => {}
-        }
-    }
-}

src/main.rs

@@ -1,3 +1,693 @@
-fn main() -> anyhow::Result<()> {
-    wgpu_template::run()
+#![feature(generic_const_exprs)]
+
+use core::sync;
+use std::cell::RefCell;
+use std::collections::HashMap;
+use std::fs::File;
+use std::rc::Rc;
+use std::sync::Arc;
+use std::sync::atomic::AtomicBool;
+use std::sync::mpsc::sync_channel;
+
+use bytemuck::Pod;
+use bytemuck::Zeroable;
+use crevice::std140::AsStd140;
+use crevice::std430::AsStd430;
+use egui::emath::fast_midpoint;
+use egui::mutex::Mutex;
+use glam::Mat4;
+use glam::Vec3;
+use glam::Vec4;
+use itertools::Itertools;
+use rand::random;
+use wgpu::BindGroup;
+use wgpu::BindGroupEntry;
+use wgpu::BindGroupLayoutDescriptor;
+use wgpu::BindGroupLayoutEntry;
+use wgpu::Buffer;
+use wgpu::BufferUsages;
+use wgpu::DepthBiasState;
+use wgpu::Device;
+use wgpu::Extent3d;
+use wgpu::Features;
+use wgpu::FragmentState;
+use wgpu::InstanceDescriptor;
+use wgpu::InstanceFlags;
+use wgpu::MemoryBudgetThresholds;
+use wgpu::NoopBackendOptions;
+use wgpu::Operations;
+use wgpu::PrimitiveState;
+use wgpu::RenderPassDepthStencilAttachment;
+use wgpu::RenderPipeline;
+use wgpu::RenderPipelineDescriptor;
+use wgpu::ShaderModuleDescriptor;
+use wgpu::ShaderStages;
+use wgpu::StencilState;
+use wgpu::Texture;
+use wgpu::TextureUsages;
+use wgpu::TextureView;
+use wgpu::VertexState;
+use wgpu::util::BufferInitDescriptor;
+use wgpu::util::DeviceExt;
+use wgpu::util::DownloadBuffer;
+use wgpu::util::StagingBelt;
+use winit::application::ApplicationHandler;
+use winit::event::DeviceEvent;
+use winit::event::WindowEvent;
+use winit::event_loop;
+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;
+
+use crate::camera::Camera;
+use crate::egui_renderer::EguiRenderer;
+use crate::producers::SineGenerator;
+use crate::voxel::cache::CacheNodeRequest;
+use crate::voxel::cache::CacheResponse;
+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;
+
+mod camera;
+mod egui_renderer;
+mod producers;
+mod voxel;
+//mod tree;
+//
+
+struct State
+{
+    instance: wgpu::Instance,
+    window: Arc<Window>,
+    device: wgpu::Device,
+    queue: wgpu::Queue,
+    size: winit::dpi::PhysicalSize<u32>,
+    surface: wgpu::Surface<'static>,
+    depth_buffer: (wgpu::Texture, wgpu::TextureView),
+    surface_format: wgpu::TextureFormat,
+    egui_renderer: EguiRenderer,
+
+    pipeline: RenderPipeline,
+    voxel_cache: Arc<Mutex<VoxelCache<4>>>,
+    insertion_debounce: bool,
+
+    camera: Camera,
+}
+
+pub unsafe fn as_raw_bytes<T: Sized>(slice: &[T]) -> &[u8]
+{
+    let size = slice.len() * size_of::<T>();
+    unsafe { std::slice::from_raw_parts(slice.as_ptr() as *const u8, size) }
+}
+
+#[derive(Clone, Copy)]
+struct Immediates
+{
+    view_proj: Mat4,
+    cam_pos: Vec3,
+    frame_timestamp: u32,
+}
+
+impl State
+{
+    async fn new(display: OwnedDisplayHandle, window: Arc<Window>) -> State
+    {
+        let instance = wgpu::Instance::new(InstanceDescriptor {
+            backends: wgpu::Backends::VULKAN,
+            display: Some(Box::new(display)),
+
+            flags: InstanceFlags::default(),
+            memory_budget_thresholds: MemoryBudgetThresholds::default(),
+            backend_options: Default::default(),
+        });
+        let adapter = instance
+            .request_adapter(&wgpu::RequestAdapterOptions::default())
+            .await
+            .unwrap();
+        let (device, queue) = adapter
+            .request_device(&wgpu::DeviceDescriptor {
+                required_features: Features::IMMEDIATES,
+                required_limits: wgpu::Limits {
+                    max_immediate_size: 112,
+                    max_storage_buffers_per_shader_stage: 16,
+                    ..Default::default()
+                },
+                ..Default::default()
+            })
+            .await
+            .unwrap();
+
+        let size = window.inner_size();
+
+        let surface = instance.create_surface(window.clone()).unwrap();
+        let cap = surface.get_capabilities(&adapter);
+        let surface_format = cap.formats[0];
+
+        let egui_renderer = EguiRenderer::new(&device, surface_format, &window);
+
+        let mut voxel_cache = VoxelCache::<4>::new(4096, device.clone(), queue.clone());
+        let id = voxel_cache.structure_table.allocate_structure(true);
+        println!("id: {id}");
+
+        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 pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+            label: Some("Voxel pipeline layout"),
+
+            bind_group_layouts: &[Some(&voxel_cache.bind_group_layout())],
+            immediate_size: size_of::<Immediates>() as u32,
+        });
+
+        let 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: &[],
+            },
+            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 state = State {
+            instance,
+            window,
+            size,
+            surface,
+            surface_format,
+            egui_renderer,
+            depth_buffer: Self::create_depth_buffer(&device, size.width, size.height),
+            queue,
+            device,
+            pipeline: chunk_pipeline,
+            voxel_cache: Arc::new(Mutex::new(voxel_cache)),
+            insertion_debounce: false,
+            camera: Default::default(),
+        };
+
+        // Configure surface for the first time
+        state.configure_surface();
+
+        state
+    }
+
+    fn get_window(&self) -> &Window
+    {
+        &self.window
+    }
+
+    fn handle_event(&mut self, event: &WindowEvent)
+    {
+        self.egui_renderer.handle_input(&self.window, event);
+        self.camera.handle_input(event);
+    }
+
+    fn handle_device_event(&mut self, event: &DeviceEvent)
+    {
+        #[allow(clippy::single_match)]
+        match event
+        {
+            winit::event::DeviceEvent::MouseMotion { delta } =>
+            {
+                self.camera.cursor_moved(delta.0 as f32, delta.1 as f32);
+            }
+
+            _ =>
+            {}
+        }
+    }
+
+    fn create_depth_buffer(device: &Device, width: u32, height: u32) -> (Texture, TextureView)
+    {
+        let texture = device.create_texture(&wgpu::wgt::TextureDescriptor {
+            label: Some("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::Depth24PlusStencil8,
+            usage: TextureUsages::RENDER_ATTACHMENT,
+            view_formats: &[wgpu::TextureFormat::Depth24PlusStencil8.add_srgb_suffix()],
+        });
+        let texture_view = texture.create_view(&wgpu::wgt::TextureViewDescriptor {
+            label: Some("depth view"),
+            ..Default::default()
+        });
+        (texture, texture_view)
+    }
+
+    fn configure_surface(&self)
+    {
+        let surface_config = wgpu::SurfaceConfiguration {
+            color_space: wgpu::SurfaceColorSpace::Auto,
+            usage: wgpu::TextureUsages::RENDER_ATTACHMENT,
+            format: self.surface_format,
+            // Request compatibility with the sRGB-format texture view we‘re going to create later.
+            view_formats: vec![self.surface_format.add_srgb_suffix()],
+            alpha_mode: wgpu::CompositeAlphaMode::Auto,
+            width: self.size.width,
+            height: self.size.height,
+            desired_maximum_frame_latency: 2,
+            present_mode: wgpu::PresentMode::AutoVsync,
+        };
+        self.surface.configure(&self.device, &surface_config);
+    }
+
+    fn resize(&mut self, new_size: winit::dpi::PhysicalSize<u32>)
+    {
+        self.size = new_size;
+        self.camera.aspect = new_size.width as f32 / new_size.height as f32;
+
+        // reconfigure the surface
+        self.configure_surface();
+        self.depth_buffer =
+            Self::create_depth_buffer(&self.device, new_size.width, new_size.height);
+    }
+
+    fn render(&mut self)
+    {
+        self.camera.update();
+
+        // Write random shit in request buffer
+        let mut belt = StagingBelt::new(self.device.clone(), size_of::<u32>() as u64);
+
+        // Create texture view.
+        // NOTE: We must handle Timeout because the surface may be unavailable
+        // (e.g., when the window is occluded on macOS).
+        let surface_texture = match self.surface.get_current_texture()
+        {
+            wgpu::CurrentSurfaceTexture::Success(texture) => texture,
+            wgpu::CurrentSurfaceTexture::Occluded | wgpu::CurrentSurfaceTexture::Timeout => return,
+            wgpu::CurrentSurfaceTexture::Suboptimal(texture) =>
+            {
+                drop(texture);
+                self.configure_surface();
+                return;
+            }
+            wgpu::CurrentSurfaceTexture::Outdated =>
+            {
+                self.configure_surface();
+                return;
+            }
+            wgpu::CurrentSurfaceTexture::Validation =>
+            {
+                unreachable!("No error scope registered, so validation errors will panic")
+            }
+            wgpu::CurrentSurfaceTexture::Lost =>
+            {
+                self.surface = self.instance.create_surface(self.window.clone()).unwrap();
+                self.configure_surface();
+                return;
+            }
+        };
+
+        let texture_view = surface_texture
+            .texture
+            .create_view(&wgpu::TextureViewDescriptor {
+                // Without add_srgb_suffix() the image we will be working with
+                // might not be "gamma correct".
+                format: Some(self.surface_format.add_srgb_suffix()),
+                ..Default::default()
+            });
+
+        // Renders a GREEN screen
+        let mut encoder = self.device.create_command_encoder(&Default::default());
+
+        {
+            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: &self.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_pipeline(&self.pipeline);
+            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(),
+            }];
+            renderpass.set_immediates(0, unsafe { as_raw_bytes(&imm) });
+            renderpass.draw(0..36, 0..1);
+
+            // End the renderpass.
+            drop(renderpass);
+        }
+
+        let requests = self.device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("dummy_dumb_dinky_aaaahhh_buffer"),
+            size: 16 * 64,
+            usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
+            mapped_at_creation: false,
+        });
+
+        self.voxel_cache
+            .lock()
+            .cache_post_render(&mut encoder, requests.clone());
+
+        // If you wanted to call any drawing commands, they would go here.
+        {
+            self.egui_renderer.begin_frame(&self.window);
+            egui::Window::new("Window ! ")
+                .resizable(true)
+                .show(self.egui_renderer.context(), |ui| ui.label("Hello !"));
+
+            let screen_descriptor = egui_wgpu::ScreenDescriptor {
+                size_in_pixels: [self.size.width, self.size.height],
+                pixels_per_point: self.window.as_ref().scale_factor() as f32,
+            };
+
+            self.egui_renderer.end_frame_and_draw(
+                &self.device,
+                &self.queue,
+                &mut encoder,
+                &self.window,
+                &texture_view,
+                screen_descriptor,
+            );
+        }
+
+        // Submit the command in the queue to execute
+        self.queue.submit([encoder.finish()]);
+        self.window.pre_present_notify();
+        self.queue.present(surface_texture);
+        if !self
+            .camera
+            .pressed_keyset
+            .contains(&winit::keyboard::KeyCode::KeyF)
+        {
+            self.insertion_debounce = false;
+        }
+
+        if (!self
+            .camera
+            .pressed_keyset
+            .contains(&winit::keyboard::KeyCode::KeyF))
+        {
+            self.insertion_debounce = true;
+            let request_count = self.voxel_cache.lock().total_request_count();
+            let cloned_cache = self.voxel_cache.clone();
+            let cloned_device = self.device.clone();
+            let cloned_queue = self.queue.clone();
+            let (tx, rx) = sync_channel(1);
+            DownloadBuffer::read_buffer(
+                &self.device.clone(),
+                &self.queue.clone(),
+                &requests.slice(0..),
+                move |buffer| {
+                    if tx.try_send(()).is_err()
+                    {
+                        return;
+                    }
+
+                    let cache_node_requests: Vec<CacheNodeRequest> =
+                        bytemuck::pod_collect_to_vec(&buffer.unwrap());
+
+                    let mut sine_gen = SineGenerator::<4>::new(4);
+                    let mut structure_nodes = vec![];
+                    let mut color_nodes = vec![];
+                    let mut location_nodes = vec![];
+                    let mut destinations = vec![];
+                    for request in cache_node_requests.iter().take(request_count as usize)
+                    {
+                        let location;
+                        let node;
+                        if request.child_index == u32::MAX
+                        {
+                            // Produce root node
+                            node = sine_gen.produce_node(0, 0, 0, 0);
+                            location = LocationPoolElement {
+                                structure_id: 0,
+                                structure_locator: NTreeNodeLocator::<4>::root().as_usize() as u32,
+                            };
+                        }
+                        else
+                        {
+                            // Figure out depth of request
+                            let locator = NTreeNodeLocator::<4>::from_usize(
+                                request.structure_locator as usize,
+                            );
+                            let depth = locator.depth();
+                            //let (x, y, z) = locator.node_location();
+
+                            let child_x = request.child_index / (4 * 4);
+                            let child_y = (request.child_index % (4 * 4)) / 4;
+                            let child_z = request.child_index % 4;
+
+                            let child_locator =
+                                locator.child(child_x as usize, child_y as usize, child_z as usize);
+
+                            let (nx, ny, nz) = child_locator.node_location();
+
+                            node = sine_gen.produce_node(depth + 1, nx, ny, nz);
+
+                            location = LocationPoolElement {
+                                structure_id: 0,
+                                structure_locator: child_locator.as_usize() as u32,
+                            };
+                        }
+
+                        structure_nodes.push(node.structure);
+                        color_nodes.push(node.colors);
+                        location_nodes.push(location);
+                        destinations.push(DestinationElement {
+                            node: request.node_index,
+                            child: request.child_index,
+                        });
+                    }
+
+                    if structure_nodes.len() != 0
+                    {
+                        let structre_node_buffer =
+                            cloned_device.create_buffer_init(&BufferInitDescriptor {
+                                label: None,
+                                contents: unsafe { as_raw_bytes(structure_nodes.as_slice()) },
+                                usage: BufferUsages::STORAGE,
+                            });
+                        let color_node_buffer =
+                            cloned_device.create_buffer_init(&BufferInitDescriptor {
+                                label: None,
+                                contents: unsafe { as_raw_bytes(color_nodes.as_slice()) },
+                                usage: BufferUsages::STORAGE,
+                            });
+                        let location_node_buffer =
+                            cloned_device.create_buffer_init(&BufferInitDescriptor {
+                                label: None,
+                                contents: unsafe { as_raw_bytes(location_nodes.as_slice()) },
+                                usage: BufferUsages::STORAGE,
+                            });
+                        let destinations_buffer =
+                            cloned_device.create_buffer_init(&BufferInitDescriptor {
+                                label: None,
+                                contents: unsafe { as_raw_bytes(destinations.as_slice()) },
+                                usage: BufferUsages::STORAGE,
+                            });
+
+                        let mut encoder = cloned_device.create_command_encoder(&Default::default());
+                        cloned_cache.lock().cache_insert(
+                            &mut encoder,
+                            &CacheResponse {
+                                structure_nodes: structre_node_buffer,
+                                color_nodes: color_node_buffer,
+                                locations: location_node_buffer,
+                                parents: destinations_buffer,
+                            },
+                        );
+                        cloned_queue.submit([encoder.finish()]);
+                    }
+                    cloned_cache.lock().next_frame();
+                },
+            );
+            println!("Total request count: {}", request_count);
+            loop
+            {
+                if rx.try_recv().is_ok()
+                {
+                    break;
+                }
+                let _ = self.device.poll(wgpu::wgt::PollType::Poll);
+            }
+            drop(rx);
+        }
+    }
+}
+
+#[derive(Default)]
+struct App
+{
+    state: Option<State>,
+}
+
+impl ApplicationHandler for App
+{
+    fn resumed(&mut self, event_loop: &ActiveEventLoop)
+    {
+        // Create window object
+        let window = Arc::new(
+            event_loop
+                .create_window(Window::default_attributes())
+                .unwrap(),
+        );
+
+        let state = pollster::block_on(State::new(
+            event_loop.owned_display_handle(),
+            window.clone(),
+        ));
+        self.state = Some(state);
+
+        window.request_redraw();
+        window.set_cursor_visible(false);
+        // window
+        //     .set_cursor_grab(winit::window::CursorGrabMode::Locked)
+        //     .unwrap();
+        let _ = window
+            .set_cursor_grab(winit::window::CursorGrabMode::Locked)
+            .or_else(|_| window.set_cursor_grab(winit::window::CursorGrabMode::Confined));
+    }
+
+    fn window_event(&mut self, event_loop: &ActiveEventLoop, _id: WindowId, event: WindowEvent)
+    {
+        let state = self.state.as_mut().unwrap();
+        state.handle_event(&event);
+        match event
+        {
+            WindowEvent::CloseRequested =>
+            {
+                println!("The close button was pressed; stopping");
+                event_loop.exit();
+            }
+            WindowEvent::RedrawRequested =>
+            {
+                state.render();
+                // Emits a new redraw requested event.
+                state.get_window().request_redraw();
+            }
+            WindowEvent::Resized(size) =>
+            {
+                // Reconfigures the size of the surface. We do not re-render
+                // here as this event is always followed up by redraw request.
+                state.resize(size);
+            }
+            _ => (),
+        }
+    }
+
+    fn device_event(
+        &mut self,
+        _event_loop: &event_loop::ActiveEventLoop,
+        _device_id: winit::event::DeviceId,
+        event: winit::event::DeviceEvent,
+    )
+    {
+        let state = self.state.as_mut().unwrap();
+        state.handle_device_event(&event);
+    }
+}
+
+fn main()
+{
+    // wgpu uses `log` for all of our logging, so we initialize a logger with the `env_logger` crate.
+    //
+    // To change the log level, set the `RUST_LOG` environment variable. See the `env_logger`
+    // documentation for more information.
+    env_logger::init();
+
+    let event_loop = EventLoop::builder().build().unwrap();
+    //let event_loop = EventLoop::builder().with_x11().build().unwrap();
+
+    // When the current loop iteration finishes, immediately begin a new
+    // iteration regardless of whether or not new events are available to
+    // process. Preferred for applications that want to render as fast as
+    // possible, like games.
+    event_loop.set_control_flow(ControlFlow::Poll);
+
+    // When the current loop iteration finishes, suspend the thread until
+    // another event arrives. Helps keeping CPU utilization low if nothing
+    // is happening, which is preferred if the application might be idling in
+    // the background.
+    // event_loop.set_control_flow(ControlFlow::Wait);
+
+    let mut app = App::default();
+    event_loop.run_app(&mut app).unwrap();
+}
+
+#[derive(AsStd140)]
+pub struct ChunkImmediate
+{
+    mvp: Mat4,
+    eye_pos: Vec3,
+    root_color: Vec4,
+    root_subdiv: bool,
+    frame_timestamp: u32,
 }

src/producers.rs

@@ -0,0 +1,88 @@
+use glam::Vec3;
+use itertools::Itertools;
+
+use crate::voxel::gpu::ExplicitNTreeNode;
+use crate::voxel::gpu::StructurePointer;
+use crate::voxel::sparse::Color;
+
+pub struct SineGenerator<const N: usize>
+{
+    chunk_power: 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
+}
+
+impl<const N: usize> SineGenerator<N>
+where
+    [(); N * N * N]:,
+{
+    pub fn new(chunk_power: usize) -> Self
+    {
+        SineGenerator { chunk_power }
+    }
+
+    pub fn produce_node(
+        &self,
+        depth: usize,
+        nx: usize,
+        ny: usize,
+        nz: usize,
+    ) -> ExplicitNTreeNode<N>
+    {
+        let node_size = N.pow((self.chunk_power - depth) as u32);
+        let child_size = node_size / N;
+        let global_size = N.pow(self.chunk_power as u32);
+
+        let mut children = vec![];
+        let mut children_color = vec![];
+
+        let gnx = nx * node_size;
+        let gny = ny * node_size;
+        let gnz = nz * node_size;
+        // Iterate on children of this node
+        for ((cx, cy), cz) in (0..N).cartesian_product(0..N).cartesian_product(0..N)
+        {
+            let gvx = gnx + cx * child_size + (child_size / 2);
+            let gvy = gny + cy * child_size + (child_size / 2);
+            let gvz = gnz + cz * child_size + (child_size / 2);
+
+            let dist = Vec3::new(gvx as f32 - 128., gvy as f32 - 128., gvz as f32 - 128.).length();
+            let child_diagonal_length = (child_size as f32 / 2.) * f32::sqrt(3.);
+
+            let alpha;
+            if (dist - 128.).abs() <= child_diagonal_length
+            {
+                children.push(StructurePointer::new(depth <= 3, false, 0));
+                alpha = if dist > 128. { 0. } else { 1. };
+            }
+            else if dist > 128.
+            {
+                children.push(StructurePointer::new(false, false, 0));
+                alpha = 0.;
+            }
+            else
+            {
+                children.push(StructurePointer::new(false, false, 0));
+                alpha = 1.;
+            }
+
+            // let alpha = 1.;
+            // children.push(StructurePointer::new(depth <= 3, false, 0));
+
+            children_color.push(Color(
+                gvx as f32 / global_size as f32,
+                gvy as f32 / global_size as f32,
+                gvz as f32 / global_size as f32,
+                alpha,
+            ));
+        }
+
+        ExplicitNTreeNode {
+            structure: std::array::from_fn(|i| children[i]),
+            colors: std::array::from_fn(|i| children_color[i]),
+        }
+    }
+}

src/state.rs

@@ -1,169 +0,0 @@
-use std::sync::Arc;
-
-use cgmath::{Matrix4, Vector3};
-use crevice::std430::AsStd430;
-use egui_wgpu::ScreenDescriptor;
-use wgpu::{Features, FeaturesWGPU, FeaturesWebGPU};
-use winit::{event::WindowEvent, window::Window};
-
-use crate::egui_renderer::EguiState;
-
-pub struct State {
-    window: Arc<Window>,
-    device: wgpu::Device,
-    queue: wgpu::Queue,
-    size: winit::dpi::PhysicalSize<u32>,
-    surface: wgpu::Surface<'static>,
-    surface_format: wgpu::TextureFormat,
-    egui_state: EguiState,
-}
-
-impl State {
-    pub async fn new(window: Arc<Window>) -> State {
-        let instance = wgpu::Instance::new(&wgpu::InstanceDescriptor::default());
-        let adapter = instance
-            .request_adapter(&wgpu::RequestAdapterOptions::default())
-            .await
-            .unwrap();
-        let (device, queue) = adapter
-            .request_device(&wgpu::DeviceDescriptor {
-                required_features: Features {
-                    features_wgpu: FeaturesWGPU::PUSH_CONSTANTS,
-                    features_webgpu: FeaturesWebGPU::empty(),
-                },
-                required_limits: wgpu::Limits {
-                    max_push_constant_size: RayMarchingPushConstants::std430_size_static() as u32,
-                    ..Default::default()
-                },
-                ..Default::default()
-            })
-            .await
-            .unwrap();
-
-        let size = window.inner_size();
-
-        let surface = instance.create_surface(window.clone()).unwrap();
-        let cap = surface.get_capabilities(&adapter);
-        let surface_format = cap.formats[0];
-
-        let state = State {
-            egui_state: EguiState::new(&device, surface_format, &window),
-
-            window,
-            device,
-            queue,
-            size,
-            surface,
-            surface_format,
-        };
-
-        // Configure surface for the first time
-        state.configure_surface();
-
-        state
-    }
-
-    pub fn get_window(&self) -> &Window {
-        &self.window
-    }
-
-    pub fn configure_surface(&self) {
-        let surface_config = wgpu::SurfaceConfiguration {
-            usage: wgpu::TextureUsages::RENDER_ATTACHMENT,
-            format: self.surface_format,
-            // Request compatibility with the sRGB-format texture view we‘re going to create later.
-            view_formats: vec![self.surface_format.add_srgb_suffix()],
-            alpha_mode: wgpu::CompositeAlphaMode::Auto,
-            width: self.size.width,
-            height: self.size.height,
-            desired_maximum_frame_latency: 2,
-            present_mode: wgpu::PresentMode::AutoVsync,
-        };
-        self.surface.configure(&self.device, &surface_config);
-    }
-
-    pub fn resize(&mut self, new_size: winit::dpi::PhysicalSize<u32>) {
-        self.size = new_size;
-
-        // reconfigure the surface
-        self.configure_surface();
-    }
-
-    pub fn handle_event(&mut self, event: &WindowEvent) {
-        self.egui_state.handle_event(&self.window, event);
-    }
-
-    pub fn render(&mut self) {
-        // Create texture view
-        let surface_texture = self
-            .surface
-            .get_current_texture()
-            .expect("failed to acquire next swapchain texture");
-        let texture_view = surface_texture
-            .texture
-            .create_view(&wgpu::TextureViewDescriptor {
-                // Without add_srgb_suffix() the image we will be working with
-                // might not be "gamma correct".
-                format: Some(self.surface_format.add_srgb_suffix()),
-                ..Default::default()
-            });
-
-        // Renders a GREEN screen
-        let mut encoder = self.device.create_command_encoder(&Default::default());
-        // Create the renderpass which will clear the screen.
-        let 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::GREEN),
-                    store: wgpu::StoreOp::Store,
-                },
-            })],
-            depth_stencil_attachment: None,
-            timestamp_writes: None,
-            occlusion_query_set: None,
-        });
-
-        // End the renderpass.
-        drop(renderpass);
-        // Egui
-        {
-            let screen_descriptor = ScreenDescriptor {
-                size_in_pixels: [self.size.width, self.size.height],
-                pixels_per_point: 1.,
-            };
-            self.egui_state.begin_frame(&self.window);
-
-            egui::Window::new("Hello Window").resizable(true).show(
-                self.egui_state.context(),
-                |ui| {
-                    ui.label("Hello, world.");
-                },
-            );
-
-            self.egui_state.end_frame_and_draw(
-                &self.device,
-                &self.queue,
-                &mut encoder,
-                &self.window,
-                &texture_view,
-                screen_descriptor,
-            );
-        }
-
-        // Submit the command in the queue to execute
-        self.queue.submit([encoder.finish()]);
-        self.window.pre_present_notify();
-        surface_texture.present();
-    }
-}
-
-#[derive(AsStd430)]
-pub struct RayMarchingPushConstants {
-    inverse_projection_matrix: Matrix4<f32>,
-    view_matrix: Matrix4<f32>,
-    camera_pos: Vector3<f32>,
-}

src/tree.rs

@@ -0,0 +1,837 @@
+use std::collections::HashMap;
+use std::collections::VecDeque;
+use std::hash::Hash;
+use std::vec;
+
+use itertools::Itertools;
+
+#[derive(Clone, Copy, PartialEq, Debug)]
+pub struct Color(pub f32, pub f32, pub f32, pub f32);
+
+pub struct NTree<const N: usize>
+{
+    structure: HashMap<NTreeLocator<N>, NTreeNode<N>>,
+    depth: u32,
+}
+
+#[derive(Clone, Copy)]
+struct NTreeNode<const N: usize>
+{
+    color: Color,
+    subdivided: bool,
+}
+
+impl<const N: usize> NTree<N>
+where
+    [(); N * N * N]:,
+{
+    pub fn constant(color: Color, depth: u32) -> Self
+    {
+        let mut structure = HashMap::new();
+        structure.insert(
+            NTreeLocator::root(),
+            NTreeNode {
+                color,
+                subdivided: false,
+            },
+        );
+        Self { structure, depth }
+    }
+
+    pub fn get(&self, x: usize, y: usize, z: usize) -> Color
+    {
+        let mut local_x = x;
+        let mut local_y = y;
+        let mut local_z = z;
+
+        let mut size = N.pow(self.depth);
+        assert!(x < size && y < size && z < size);
+
+        let mut current_loc = NTreeLocator::<N>::root();
+        let mut current_node = *self.structure.get(&current_loc).unwrap();
+
+        loop
+        {
+            if !current_node.subdivided
+            {
+                return current_node.color;
+            }
+
+            size /= N;
+
+            // Descend
+            let child_x = local_x / size;
+            let child_y = local_y / size;
+            let child_z = local_z / size;
+
+            // Node is subdivided, descend
+            current_loc = current_loc.get_child(child_x, child_y, child_z);
+            current_node = *self.structure.get(&current_loc).unwrap();
+
+            local_x -= child_x * size;
+            local_y -= child_y * size;
+            local_z -= child_z * size;
+        }
+    }
+
+    pub fn from_arrays_old(chunk: &[Color], depth: u32) -> Self
+    {
+        let chunk_width = N.pow(depth);
+        let mut structure = HashMap::with_capacity(chunk_width * chunk_width * chunk_width);
+        println!("Inserting voxels");
+        for ((x, y), z) in (0..chunk_width)
+            .cartesian_product(0..chunk_width)
+            .cartesian_product(0..chunk_width)
+        {
+            structure.insert(
+                NTreeLocator::<N>::from_depth_coords(x, y, z, depth as usize),
+                NTreeNode::<N> {
+                    color: chunk[x + y * chunk_width + z * chunk_width * chunk_width],
+                    subdivided: false,
+                },
+            );
+        }
+
+        println!("Starting bottom up");
+        let mut current_size = 1;
+        for d in (0..depth).rev()
+        {
+            println!("depth: {d}");
+            current_size *= N;
+            for ((x, y), z) in (0..(chunk_width / current_size))
+                .cartesian_product(0..(chunk_width / current_size))
+                .cartesian_product(0..(chunk_width / current_size))
+            {
+                let loc = NTreeLocator::<N>::from_depth_coords(x, y, z, d as usize);
+                let children = (0..N)
+                    .cartesian_product(0..N)
+                    .cartesian_product(0..N)
+                    .map(|((cx, cy), cz)| structure.get(&loc.get_child(cx, cy, cz)).unwrap())
+                    .collect::<Vec<_>>();
+
+                let (can_merge, _) = children.iter().fold((true, None), |acc, child| match acc
+                {
+                    (_, None) => (true, Some(child.color)),
+                    (b, Some(color)) =>
+                    {
+                        (b && !child.subdivided && color == child.color, Some(color))
+                    }
+                });
+
+                let node;
+                if can_merge
+                {
+                    node = NTreeNode::<N> {
+                        color: children[0].color,
+                        subdivided: false,
+                    };
+
+                    //Remove merged childrenn
+                    (0..N)
+                        .cartesian_product(0..N)
+                        .cartesian_product(0..N)
+                        .for_each(|((cx, cy), cz)| {
+                            structure.remove(&loc.get_child(cx, cy, cz)).unwrap();
+                        });
+                }
+                else
+                {
+                    node = NTreeNode::<N> {
+                        color: Color::average(
+                            children
+                                .iter()
+                                .map(|x| x.color)
+                                .collect::<Vec<_>>()
+                                .as_slice(),
+                        ),
+                        subdivided: true,
+                    };
+                }
+                structure.insert(loc, node);
+            }
+        }
+
+        Self { structure, depth }
+    }
+
+    pub fn from_arrays(chunk: &[Color], depth: usize) -> Self
+    {
+        let width = N.pow(depth as u32);
+        let mut structure = HashMap::new();
+        let mut taken = vec![0; chunk.len()]; // Whether or not the current voxel has been added
+
+        // Voxel insertion/combination pass
+        for (y, z) in (0..width).cartesian_product(0..width)
+        {
+            let mut x = 0;
+            while x < width
+            {
+                if taken[x + y * width + z * width * width] != 0
+                {
+                    x += taken[x + y * width + z * width * width];
+                    continue;
+                }
+
+                let max_combination_depth = trailing_zeroes::<N>(x)
+                    .unwrap_or(depth)
+                    .min(trailing_zeroes::<N>(y).unwrap_or(depth))
+                    .min(trailing_zeroes::<N>(z).unwrap_or(depth));
+
+                let prev_color = chunk[x + y * width + z * width * width];
+                let mut locator = NTreeLocator::<N>::from_depth_coords(x, y, z, depth);
+                //let mut insertion_depth = depth;
+                let mut block_width = 1;
+                'depth_loop: for depth in 1..=max_combination_depth
+                {
+                    let combination_width = N.pow(depth as u32);
+                    for ((sx, sy), sz) in (0..combination_width)
+                        .cartesian_product(0..combination_width)
+                        .cartesian_product(0..combination_width)
+                    {
+                        let voxel_color =
+                            chunk[(x + sx) + (y + sy) * width + (z + sz) * width * width];
+                        let voxel_taken =
+                            taken[(x + sx) + (y + sy) * width + (z + sz) * width * width];
+                        if prev_color != voxel_color || voxel_taken != 0
+                        {
+                            // Cannot merge further
+                            break 'depth_loop;
+                        }
+                    }
+
+                    // At this point, voxel in combination_width^3 block can be merged
+                    block_width = combination_width;
+                    //insertion_depth -= 1;
+                    locator = locator.get_parent();
+                }
+
+                structure.insert(
+                    locator,
+                    NTreeNode {
+                        color: prev_color,
+                        subdivided: false,
+                    },
+                );
+
+                // Mark as taken
+                for ((sx, sy), sz) in (0..block_width)
+                    .cartesian_product(0..block_width)
+                    .cartesian_product(0..block_width)
+                {
+                    taken[(x + sx) + (y + sy) * width + (z + sz) * width * width] = block_width;
+                }
+
+                x += block_width;
+            }
+        }
+
+        // Increase pass
+        for d in (0..=(depth - 1)).rev()
+        {
+            // Iterate on blocks of depth d
+            let block_count = N.pow(d as u32); // Number of such blocks along axis
+            let block_width = width / block_count; // Number of such blocks along axis
+            for ((bx, by), bz) in (0..(block_count))
+                .cartesian_product(0..(block_count))
+                .cartesian_product(0..(block_count))
+            {
+                // Get how was the origin voxel merged
+                let merged_width = taken[(bx * block_width)
+                    + (by * block_width) * width
+                    + (bz * block_width) * width * width];
+
+                if merged_width >= block_width
+                {
+                    // Current voxels have been merged in bigger or equal block
+                    continue;
+                }
+
+                let locator = NTreeLocator::<N>::from_depth_coords(bx, by, bz, d);
+
+                // Otherwise, merge
+
+                // Children CANNOT be merged, Otherwise they would have been already
+                // Compute average color
+
+                let colors = locator
+                    .iter_children()
+                    .map(|loc| structure.get(&loc).unwrap().color)
+                    .collect::<Vec<_>>();
+                structure.insert(
+                    locator,
+                    NTreeNode::<N> {
+                        color: Color::average(&colors),
+                        subdivided: true,
+                    },
+                );
+            }
+        }
+
+        Self {
+            structure,
+            depth: depth as u32,
+        }
+    }
+
+    pub fn set(&mut self, x: usize, y: usize, z: usize, color: Color)
+    {
+        let mut local_x = x;
+        let mut local_y = y;
+        let mut local_z = z;
+
+        let mut size = N.pow(self.depth);
+        assert!(x < size && y < size && z < size);
+
+        let mut current_loc = NTreeLocator::<N>::root();
+        let mut current_node = *self.structure.get(&current_loc).unwrap();
+        let mut current_depth = 0;
+
+        loop
+        {
+            // Subnodes alread with correct color
+            if current_node.color == color
+            {
+                return;
+            }
+
+            if current_depth == self.depth
+            {
+                current_node.color = color;
+                self.structure.get_mut(&current_loc).unwrap().color = color;
+                break;
+            }
+
+            if !current_node.subdivided
+            {
+                // Have to subdivide
+                let sub_child = NTreeNode::<N> {
+                    color: current_node.color,
+                    subdivided: false,
+                };
+                // Set node as subdivided
+                self.structure.get_mut(&current_loc).unwrap().subdivided = true;
+                current_node.subdivided = true;
+
+                // Insert new children
+                current_loc.iter_children().for_each(|x| {
+                    self.structure.insert(x, sub_child);
+                });
+            }
+
+            if current_node.subdivided
+            {
+                size /= N;
+                let child_x = local_x / size;
+                let child_y = local_y / size;
+                let child_z = local_z / size;
+
+                // Node is subdivided, descend
+                current_loc = current_loc.get_child(child_x, child_y, child_z);
+                current_node = *self.structure.get(&current_loc).unwrap();
+                current_depth += 1;
+
+                local_x -= child_x * size;
+                local_y -= child_y * size;
+                local_z -= child_z * size;
+            }
+        }
+
+        // Insertion has been done, travel back up to optimise
+        loop
+        {
+            // Try to simplify, compute average
+            if current_node.subdivided
+            {
+                let (compressable, color) = current_loc
+                    .iter_children()
+                    .map(|x| self.structure.get(&x).unwrap())
+                    .fold((true, None), |state, child| match state
+                    {
+                        (_, None) => (!child.subdivided, Some(child.color)),
+                        (b, Some(c)) => (b && !child.subdivided && c == child.color, Some(c)),
+                    });
+
+                if compressable
+                {
+                    // Compress
+                    *self.structure.get_mut(&current_loc).unwrap() = NTreeNode {
+                        color: color.unwrap(),
+                        subdivided: false,
+                    };
+
+                    // Remove children
+                    current_loc.iter_children().for_each(|x| {
+                        self.structure.remove(&x).unwrap();
+                    });
+                }
+                else
+                {
+                    // Update average color
+                    let colors = current_loc
+                        .iter_children()
+                        .map(|x| self.structure.get(&x).unwrap().color)
+                        .collect::<Vec<_>>();
+                    self.structure.get_mut(&current_loc).unwrap().color =
+                        Color::average(colors.as_slice());
+                }
+            }
+
+            // Travel up
+            if current_depth == 0
+            {
+                break; // Finished
+            }
+
+            current_loc = current_loc.get_parent();
+            current_node = *self.structure.get(&current_loc).unwrap();
+            current_depth -= 1;
+        }
+    }
+
+    pub fn to_gpu_rep(&self) -> (Color, Vec<GPUStructureTile<N>>, Vec<[Color; N * N * N]>)
+    {
+        // No root + group by child group
+        let tile_count = (self.structure.len() - 1) / (N * N * N);
+        if tile_count == 0
+        {
+            return (
+                self.structure
+                    .get(&NTreeLocator::<N>::root())
+                    .unwrap()
+                    .color,
+                vec![],
+                vec![],
+            );
+        }
+
+        let mut structure_tiles = vec![GPUStructureTile::<N>::zero(); tile_count];
+        let mut color_tiles = vec![[Color(0., 0., 0., 0.); N * N * N]; tile_count];
+
+        let mut current_tile = 0usize;
+        let mut queue = VecDeque::new();
+        queue.push_back((NTreeLocator::<N>::root(), current_tile));
+        current_tile += 1;
+        while !queue.is_empty()
+        {
+            let (loc, dest_tile) = queue.pop_front().unwrap();
+
+            let children = loc
+                .iter_children()
+                .map(|x| (x, self.structure.get(&x).unwrap()))
+                .collect::<Vec<_>>();
+
+            let mut structure_tile = GPUStructureTile::<N>::zero();
+            let mut color_tile = [Color(0., 0., 0., 0.); N * N * N];
+
+            for (i, (child_loc, child)) in children.into_iter().enumerate()
+            {
+                structure_tile.children[i] = GPUStructureTileIndex::new(
+                    child.subdivided,
+                    child.subdivided,
+                    current_tile as u32,
+                );
+                color_tile[i] = child.color;
+
+                if child.subdivided
+                {
+                    queue.push_back((child_loc, current_tile));
+                    current_tile += 1;
+                }
+            }
+
+            structure_tiles[dest_tile] = structure_tile;
+            color_tiles[dest_tile] = color_tile;
+        }
+
+        println!(
+            "node count: {}, current_tile: {}",
+            self.structure.len(),
+            current_tile
+        );
+        (
+            self.structure
+                .get(&NTreeLocator::<N>::root())
+                .unwrap()
+                .color,
+            structure_tiles,
+            color_tiles,
+        )
+    }
+}
+
+#[derive(Clone, Copy)]
+pub struct GPUStructureTile<const N: usize>
+where
+    [(); N * N * N]:,
+{
+    children: [GPUStructureTileIndex; N * N * N],
+}
+
+impl<const N: usize> GPUStructureTile<N>
+where
+    [(); N * N * N]:,
+{
+    pub fn zero() -> Self
+    {
+        GPUStructureTile {
+            children: [GPUStructureTileIndex::zero(); N * N * N],
+        }
+    }
+}
+
+#[derive(Clone, Copy)]
+pub struct GPUStructureTileIndex(u32);
+
+impl GPUStructureTileIndex
+{
+    pub fn zero() -> Self
+    {
+        GPUStructureTileIndex(0)
+    }
+
+    pub fn new(subdivided: bool, leaf: bool, ptr: u32) -> Self
+    {
+        assert_eq!((ptr >> 30), 0);
+        GPUStructureTileIndex(ptr | ((leaf as u32) << 30) | ((subdivided as u32) << 31))
+    }
+
+    pub fn index(&self) -> u32
+    {
+        self.0 & 0x3FFFFFFFu32
+    }
+
+    pub fn subdivided(&self) -> bool
+    {
+        (self.0 >> 31) == 1
+    }
+
+    pub fn leaf(&self) -> bool
+    {
+        ((self.0 >> 30) & 1) == 1
+    }
+}
+
+impl Color
+{
+    pub fn average(colors: &[Color]) -> Color
+    {
+        let sum = colors
+            .iter()
+            .copied()
+            .reduce(|Color(r_a, g_a, b_a, a_a), Color(r_b, g_b, b_b, a_b)| {
+                Color(r_a + r_b, g_a + g_b, b_a + b_b, a_a + a_b)
+            })
+            .unwrap_or(Color(0., 0., 0., 0.));
+        //let len = colors.len().max(1) as f32;
+        let len = colors.iter().map(|x| x.3).sum::<f32>();
+
+        Color(sum.0 / len, sum.1 / len, sum.2 / len, sum.3 / len)
+    }
+}
+
+#[derive(Clone, Copy, Hash, PartialEq, Eq, Debug)]
+pub struct NTreeLocator<const N: usize>(usize, usize, usize);
+
+impl<const N: usize> NTreeLocator<N>
+{
+    pub fn root() -> Self
+    {
+        Self(1, 1, 1)
+    }
+
+    pub fn get_child(&self, child_x: usize, child_y: usize, child_z: usize) -> Self
+    {
+        assert!(child_x < N && child_y < N && child_z < N);
+        let mut new_loc_x = self.0;
+        let mut new_loc_y = self.1;
+        let mut new_loc_z = self.2;
+
+        // Shift to left three times
+        new_loc_x *= N;
+        new_loc_x += child_x;
+
+        new_loc_y *= N;
+        new_loc_y += child_y;
+
+        new_loc_z *= N;
+        new_loc_z += child_z;
+
+        Self(new_loc_x, new_loc_y, new_loc_z)
+    }
+
+    pub fn from_depth_coords(x: usize, y: usize, z: usize, depth: usize) -> Self
+    {
+        if depth == 0
+        {
+            return Self::root();
+        }
+        let off = N.pow(depth as u32);
+        Self(off + x, off + y, off + z)
+    }
+
+    pub fn iter_children(&self) -> impl Iterator<Item = NTreeLocator<N>>
+    {
+        (0..N)
+            .cartesian_product(0..N)
+            .cartesian_product(0..N)
+            .map(|((x, y), z)| self.get_child(x, y, z))
+    }
+
+    pub fn get_parent(&self) -> Self
+    {
+        if self.0 == 0
+        {
+            *self
+        }
+        else
+        {
+            Self(self.0 / N, self.1 / N, self.2 / N)
+        }
+    }
+
+    pub fn get_local_location(&self) -> (usize, usize, usize)
+    {
+        let mut loc = self.0;
+        let z = loc % N;
+        loc /= N;
+
+        let y = loc % N;
+        loc /= N;
+
+        let x = loc % N;
+
+        (x, y, z)
+    }
+
+    pub fn is_root(&self) -> bool
+    {
+        self.0 == 1
+    }
+}
+
+fn trailing_zeroes<const N: usize>(mut n: usize) -> Option<usize>
+{
+    if n == 0
+    {
+        return None;
+    }
+    let mut i = 0;
+    while n.is_multiple_of(N)
+    // n % N == 0
+    {
+        i += 1;
+        n /= N;
+    }
+    Some(i)
+}
+
+#[cfg(test)]
+mod test
+{
+
+    use crate::voxel::trailing_zeroes;
+    use itertools::Itertools;
+    use rand::Rng;
+    use rand::RngExt;
+
+    use crate::voxel::Color;
+    use crate::voxel::NTree;
+    use crate::voxel::NTreeLocator;
+
+    #[test]
+    pub fn constant()
+    {
+        let color = Color(0.5, 0.3, 0.5, 1.);
+        let depth = 5;
+        const N: usize = 3;
+        let width = N.pow(depth);
+        let ntree = NTree::<N>::constant(color, depth);
+
+        for ((x, y), z) in (0..width)
+            .cartesian_product(0..width)
+            .cartesian_product(0..width)
+        {
+            assert_eq!(ntree.get(x, y, z), color);
+        }
+    }
+
+    #[test]
+    pub fn full_insert()
+    {
+        const DEPTH: u32 = 4;
+        const N: usize = 3;
+        const WIDTH: usize = N.pow(DEPTH);
+        let mut ntree = NTree::<N>::constant(Color(1., 0., 0., 0.), DEPTH);
+
+        let mut rng = rand::rng();
+        let mut storage = vec![vec![vec![Color(0., 0., 0., 0.); WIDTH]; WIDTH]; WIDTH];
+        for ((x, y), z) in (0..WIDTH)
+            .cartesian_product(0..WIDTH)
+            .cartesian_product(0..WIDTH)
+        {
+            let color = Color(rng.random(), rng.random(), rng.random(), rng.random());
+            storage[x][y][z] = color;
+            ntree.set(x, y, z, color);
+        }
+
+        for ((x, y), z) in (0..WIDTH)
+            .cartesian_product(0..WIDTH)
+            .cartesian_product(0..WIDTH)
+        {
+            let color = ntree.get(x, y, z);
+            assert_eq!(storage[x][y][z], color);
+        }
+        println!("Total nodes {}", ntree.structure.len());
+    }
+
+    #[test]
+    pub fn full_insert_bottom_up()
+    {
+        const DEPTH: u32 = 4;
+        const N: usize = 3;
+        const WIDTH: usize = N.pow(DEPTH);
+
+        let mut rng = rand::rng();
+        let mut storage = vec![Color(0., 0., 0., 0.); WIDTH * WIDTH * WIDTH];
+        for ((x, y), z) in (0..WIDTH)
+            .cartesian_product(0..WIDTH)
+            .cartesian_product(0..WIDTH)
+        {
+            let color = Color(rng.random(), rng.random(), rng.random(), rng.random());
+            storage[x + WIDTH * y + WIDTH * WIDTH * z] = color;
+        }
+        let ntree = NTree::<N>::from_arrays(&storage, DEPTH as usize);
+
+        for ((x, y), z) in (0..WIDTH)
+            .cartesian_product(0..WIDTH)
+            .cartesian_product(0..WIDTH)
+        {
+            let color = ntree.get(x, y, z);
+            assert_eq!(storage[x + WIDTH * y + WIDTH * WIDTH * z], color);
+        }
+        println!("Total nodes {}", ntree.structure.len());
+    }
+
+    #[test]
+    pub fn full_insert_const()
+    {
+        const DEPTH: u32 = 4;
+        const N: usize = 3;
+        const WIDTH: usize = N.pow(DEPTH);
+        const NEW_COLOR: Color = Color(0., 1., 0., 1.);
+        let mut ntree = NTree::<N>::constant(Color(1., 0., 1., 0.), DEPTH);
+
+        for ((x, y), z) in (0..WIDTH)
+            .cartesian_product(0..WIDTH)
+            .cartesian_product(0..WIDTH)
+        {
+            ntree.set(x, y, z, NEW_COLOR);
+        }
+
+        for ((x, y), z) in (0..WIDTH)
+            .cartesian_product(0..WIDTH)
+            .cartesian_product(0..WIDTH)
+        {
+            let color = ntree.get(x, y, z);
+            assert_eq!(NEW_COLOR, color);
+        }
+
+        println!("Total nodes {}", ntree.structure.len());
+    }
+
+    #[test]
+    pub fn full_insert_quantized()
+    {
+        const DEPTH: u32 = 4;
+        const N: usize = 3;
+        const WIDTH: usize = N.pow(DEPTH);
+        let mut ntree = NTree::<N>::constant(Color(0., 0., 0., 0.), DEPTH);
+
+        let mut rng = rand::rng();
+        let mut storage = vec![vec![vec![Color(0., 0., 0., 0.); WIDTH]; WIDTH]; WIDTH];
+        for ((x, y), z) in (0..WIDTH)
+            .cartesian_product(0..WIDTH)
+            .cartesian_product(0..WIDTH)
+        {
+            let nbr = rng.random::<f32>();
+            let mut color = Color(0., 1., 0., 0.);
+            // Only 1 percent of different color
+            if nbr > 0.01
+            {
+                color = Color(1., 0., 0., 0.);
+            }
+            storage[x][y][z] = color;
+            ntree.set(x, y, z, color);
+        }
+
+        for ((x, y), z) in (0..WIDTH)
+            .cartesian_product(0..WIDTH)
+            .cartesian_product(0..WIDTH)
+        {
+            let color = ntree.get(x, y, z);
+            assert_eq!(storage[x][y][z], color);
+        }
+        println!("Total nodes {}", ntree.structure.len());
+    }
+
+    #[test]
+    pub fn gpu_rep_quantized()
+    {
+        const DEPTH: u32 = 4;
+        const N: usize = 3;
+        const WIDTH: usize = N.pow(DEPTH);
+        let mut ntree = NTree::<N>::constant(Color(0., 0., 0., 0.), DEPTH);
+
+        let mut rng = rand::rng();
+        let mut storage = vec![vec![vec![Color(0., 0., 0., 0.); WIDTH]; WIDTH]; WIDTH];
+        for ((x, y), z) in (0..WIDTH)
+            .cartesian_product(0..WIDTH)
+            .cartesian_product(0..WIDTH)
+        {
+            let nbr = rng.random::<f32>();
+            let mut color = Color(0., 1., 0., 0.);
+            // Only 1 percent of different color
+            if nbr > 0.01
+            {
+                color = Color(1., 0., 0., 0.);
+            }
+            storage[x][y][z] = color;
+            ntree.set(x, y, z, color);
+        }
+
+        let (_color, _a, _b) = ntree.to_gpu_rep();
+
+        drop(_a);
+        drop(_b);
+    }
+
+    #[test]
+    pub fn tree_locator_from_depth_coords()
+    {
+        const N: usize = 4;
+        assert_eq!(
+            NTreeLocator::<N>::root(),
+            NTreeLocator::<N>::from_depth_coords(1, 2, 3, 0)
+        );
+
+        assert_eq!(
+            NTreeLocator::<N>::root().get_child(1, 2, 3),
+            NTreeLocator::<N>::from_depth_coords(1, 2, 3, 1)
+        );
+
+        assert_eq!(
+            NTreeLocator::<N>::root()
+                .get_child(1, 1, 1)
+                .get_child(1, 2, 3),
+            NTreeLocator::<N>::from_depth_coords(4 + 1, 4 + 2, 4 + 3, 2)
+        );
+    }
+
+    #[test]
+    fn trailing_zeroes_base10()
+    {
+        let n = 2139800000;
+        assert_eq!(trailing_zeroes::<10>(n), Some(5));
+        assert_eq!(trailing_zeroes::<10>(0), None);
+    }
+}

src/voxel.rs

@@ -0,0 +1,4 @@
+pub mod cache;
+pub mod gpu;
+pub mod pipeline;
+pub mod sparse;

src/voxel/gpu.rs

@@ -0,0 +1,27 @@
+use bytemuck::Pod;
+use bytemuck::Zeroable;
+
+use crate::voxel::sparse::Color;
+
+#[derive(Clone, Copy, Pod, Zeroable)]
+#[repr(transparent)]
+pub struct StructurePointer(pub u32);
+
+impl StructurePointer
+{
+    pub fn new(subdivided: bool, pointer_valid: bool, pointer: u32) -> Self
+    {
+        assert!(pointer >> 30 == 0);
+        StructurePointer((subdivided as u32) << 31 | (pointer_valid as u32) << 30 | pointer)
+    }
+}
+
+#[derive(Clone, Copy, Zeroable)]
+#[repr(C)]
+pub struct ExplicitNTreeNode<const N: usize>
+where
+    [(); N * N * N]:,
+{
+    pub structure: [StructurePointer; N * N * N],
+    pub colors: [Color; N * N * N],
+}

src/voxel/pipeline.rs

@@ -0,0 +1,799 @@
+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/sparse.rs

@@ -0,0 +1,525 @@
+use std::collections::HashMap;
+use std::collections::VecDeque;
+
+use bytemuck::Pod;
+use bytemuck::Zeroable;
+use itertools::Itertools;
+
+use crate::voxel::gpu::ExplicitNTreeNode;
+use crate::voxel::gpu::StructurePointer;
+
+#[derive(Debug, Clone, Copy, PartialEq, Pod, Zeroable)]
+#[repr(C)]
+pub struct Color(pub f32, pub f32, pub f32, pub f32);
+
+impl Color
+{
+    pub fn average<I: Iterator<Item = Color>>(iter: I) -> Self
+    {
+        let mut c = Color(0., 0., 0., 0.);
+        let mut count = 0;
+        for Color(r, g, b, a) in iter
+        {
+            count += 1;
+            c.0 += r;
+            c.1 += g;
+            c.2 += b;
+            c.3 += a;
+        }
+        c.0 /= count as f32;
+        c.1 /= count as f32;
+        c.2 /= count as f32;
+        c.3 /= count as f32;
+        c
+    }
+}
+
+#[derive(Debug, Clone)]
+pub struct NTree<const N: usize>
+{
+    map: HashMap<NTreeNodeLocator<N>, NTreeNode>,
+
+    // Depth is the parameter such that
+    // N^depth == max_number_of_voxels
+    // one root => depth = 1
+    depth: usize,
+}
+
+pub struct NTreeIter<'a, const N: usize>
+{
+    tree: &'a NTree<N>,
+    dfs_stack: Vec<NTreeNodeLocator<N>>,
+}
+
+#[derive(Debug, Clone, Copy)]
+pub struct NTreeNode
+{
+    color: Color,
+    subdivided: bool,
+}
+
+#[derive(Debug, Clone, Copy, Hash, PartialEq, Eq)]
+pub struct NTreeNodeLocator<const N: usize>(usize);
+
+impl<const N: usize> NTree<N>
+where
+    [(); N * N * N]:,
+{
+    pub fn constant(color: Color, depth: usize) -> Self
+    {
+        let mut map = HashMap::new();
+        let root = NTreeNode {
+            color,
+            subdivided: false,
+        };
+        map.insert(NTreeNodeLocator::<N>::root(), root);
+        Self { map, depth }
+    }
+
+    pub fn set(&mut self, x: usize, y: usize, z: usize, color: Color)
+    {
+        self.set_at_depth(x, y, z, self.depth, color);
+    }
+
+    pub fn get(&self, x: usize, y: usize, z: usize) -> Color
+    {
+        assert!(x < N.pow(self.depth as u32));
+        assert!(y < N.pow(self.depth as u32));
+        assert!(z < N.pow(self.depth as u32));
+
+        // Try to walk down the tree, and create necessary nodes
+        let mut current_node_loc = NTreeNodeLocator::<N>::root();
+        let mut current_node_size = N.pow(self.depth as u32);
+        let mut current_node = self.map.get(&current_node_loc).unwrap();
+
+        loop
+        {
+            if !current_node.subdivided
+            {
+                return current_node.color;
+            }
+
+            // Descend to relevant node,
+            let child_node_size = current_node_size / N;
+            let child_x = (x % current_node_size) / child_node_size;
+            let child_y = (y % current_node_size) / child_node_size;
+            let child_z = (z % current_node_size) / child_node_size;
+            current_node_size = child_node_size;
+
+            current_node_loc = current_node_loc.child(child_x, child_y, child_z);
+            current_node = self.map.get(&current_node_loc).unwrap();
+        }
+    }
+
+    pub fn set_at_depth(&mut self, x: usize, y: usize, z: usize, depth: usize, color: Color)
+    {
+        assert!(x < N.pow(depth as u32));
+        assert!(y < N.pow(depth as u32));
+        assert!(z < N.pow(depth as u32));
+
+        // Try to walk down the tree, and create necessary nodes
+        let mut current_node_loc = NTreeNodeLocator::<N>::root();
+        let mut current_node_size = N.pow(self.depth as u32);
+        let mut current_depth = 0;
+        let mut current_node = self.map.get_mut(&current_node_loc).unwrap();
+
+        while current_depth != depth
+        {
+            if current_node.color == color
+            {
+                // Color is already correct, work is finished
+                return;
+            }
+
+            // Current is not subdivided, and color is different => subdivide
+            if !current_node.subdivided
+            {
+                current_node.subdivided = true;
+                let new_nodes = NTreeNode {
+                    color: current_node.color,
+                    subdivided: false,
+                };
+                current_node_loc.children().into_iter().for_each(|child| {
+                    self.map.insert(child, new_nodes);
+                });
+            }
+
+            // Descend to relevant node,
+            let child_node_size = current_node_size / N;
+            let child_x = (x % current_node_size) / child_node_size;
+            let child_y = (y % current_node_size) / child_node_size;
+            let child_z = (z % current_node_size) / child_node_size;
+            current_node_size = child_node_size;
+            current_depth += 1;
+
+            current_node_loc = current_node_loc.child(child_x, child_y, child_z);
+            current_node = self.map.get_mut(&current_node_loc).unwrap();
+        }
+
+        if current_node.subdivided
+        {
+            self.make_constant(current_node_loc, color);
+        }
+        else
+        {
+            current_node.color = color;
+        }
+
+        // We need to go back up the hierarchy to combine nodes
+        // current_node is the modified node, go one up
+
+        loop
+        {
+            // Go one node up
+            current_depth -= 1;
+            current_node_loc = current_node_loc.parent().unwrap();
+            //current_node_size *= N;
+
+            // Access all children's colors
+            let children_color_avg = current_node_loc
+                .children()
+                .iter()
+                .map(|child_loc| self.map.get(child_loc).unwrap().color)
+                // AGGREGATE AND AVERAGE HERE !!
+                .next()
+                .unwrap();
+
+            current_node = self.map.get_mut(&current_node_loc).unwrap();
+            current_node.color = children_color_avg;
+
+            if current_depth == 0
+            {
+                break;
+            }
+        }
+    }
+
+    fn make_constant(&mut self, locator: NTreeNodeLocator<N>, color: Color)
+    {
+        let mut dfs_stack = Vec::new();
+        dfs_stack.push(locator);
+
+        while let Some(current) = dfs_stack.pop()
+        {
+            // Remove element
+            let removed = self.map.remove(&current);
+            let handle_children = removed.is_some_and(|x| x.subdivided);
+
+            // If children have to be handled, push on stack
+            if handle_children
+            {
+                dfs_stack.extend(current.children());
+            }
+        }
+
+        self.map.insert(
+            locator,
+            NTreeNode {
+                color,
+                subdivided: false,
+            },
+        );
+    }
+
+    pub fn from_array(array: &[Color], depth: usize) -> Self
+    {
+        let mut structure = HashMap::new();
+        // Per voxel in the array :
+        // 0 -> Voxel not merged yet
+        // _ -> size of the block in which the voxel was merged
+        let mut merged_size = vec![0; array.len()];
+        let width = N.pow(depth as u32);
+
+        for (x, y) in (0..width).cartesian_product(0..width)
+        {
+            let mut z = 0;
+            while z < width
+            {
+                if merged_size[x + y * width + z * width * width] != 0
+                {
+                    // If current voxel has already been accounted for -> skip
+                    // We stumbled on a big merged block, we can skip it as a whole on the inner
+                    // dimension
+                    z += merged_size[x + y * width + z * width * width];
+                    continue;
+                }
+
+                // Check smallest depth aligned with current position
+                let x_trailing = base_trailing_zeroes::<N>(x).unwrap_or(depth as u32);
+                let y_trailing = base_trailing_zeroes::<N>(y).unwrap_or(depth as u32);
+                let z_trailing = base_trailing_zeroes::<N>(z).unwrap_or(depth as u32);
+
+                // Biggest block size power of which this voxel can be a lower corner
+                let block_power = x_trailing.min(y_trailing.min(z_trailing));
+
+                // Try to merge for each depth
+                let mut max_power_merge = 0; // Single voxel can always be merged
+                let color = array[x + y * width + z * width * width];
+                let mut locator = NTreeNodeLocator::<N>::new(depth, x, y, z);
+                'power_search: for power in 1..=block_power
+                {
+                    let block_size = N.pow(power);
+                    // Check if candidate block has the same color
+                    for ((merge_x, merge_y), merge_z) in (0..block_size)
+                        .cartesian_product(0..block_size)
+                        .cartesian_product(0..block_size)
+                    {
+                        if array
+                            [(x + merge_x) + (y + merge_y) * width + (z + merge_z) * width * width]
+                            != color
+                        {
+                            break 'power_search;
+                        }
+                    }
+
+                    // Update biggest candidate block
+                    locator = locator.parent().unwrap();
+                    max_power_merge = power;
+                }
+
+                // We know which size we can know merge
+                // Say we accounted for those voxels :
+                let block_size = N.pow(max_power_merge);
+                for ((merge_x, merge_y), merge_z) in (0..block_size)
+                    .cartesian_product(0..block_size)
+                    .cartesian_product(0..block_size)
+                {
+                    merged_size
+                        [(x + merge_x) + (y + merge_y) * width + (z + merge_z) * width * width] =
+                        block_size;
+                }
+
+                // Add new found block in our structure
+                // If power is depth, this is the whole structure, thus depth is 0
+                let node = NTreeNode {
+                    color,
+                    subdivided: false,
+                };
+                structure.insert(locator, node);
+                z += block_size;
+            }
+        }
+
+        // The structure now consists of the leaf most nodes, we need to rebuild the hierarchy
+        // bottom up
+        for current_power in 1..=depth
+        {
+            // Iterate on the blocks of this size
+            let block_size = N.pow(current_power as u32);
+            let block_count = width / block_size;
+
+            for ((block_x, block_y), block_z) in (0..block_count)
+                .cartesian_product(0..block_count)
+                .cartesian_product(0..block_count)
+            {
+                // Check the lower corner of this block
+                let corner = merged_size[(block_x * block_size)
+                    + (block_y * block_size) * width
+                    + (block_z * block_size) * width * width];
+                if corner >= block_size
+                {
+                    // Children are merged in bigger block
+                    continue;
+                }
+
+                // Add hierarchy node
+                let locator =
+                    NTreeNodeLocator::<N>::new(depth - current_power, block_x, block_y, block_z);
+
+                // Get average color of children
+                let color = Color::average(
+                    locator
+                        .children()
+                        .iter()
+                        .map(|child_loc| structure.get(child_loc).unwrap().color),
+                );
+
+                let hierarchy_node = NTreeNode {
+                    color,
+                    subdivided: true,
+                };
+                structure.insert(locator, hierarchy_node);
+            }
+        }
+
+        Self {
+            map: structure,
+            depth,
+        }
+    }
+
+    pub fn as_explicit_structure(&self) -> (Color, Vec<ExplicitNTreeNode<N>>)
+    {
+        let root = self.map.get(&NTreeNodeLocator::root()).unwrap();
+        if !root.subdivided
+        {
+            return (root.color, vec![]);
+        }
+
+        let new_node = ExplicitNTreeNode {
+            structure: [StructurePointer(0); _],
+            colors: [Color(0., 0., 0., 0.); _],
+        };
+        let mut explicit_structure = vec![new_node.clone()];
+        let mut bfs_queue = VecDeque::new();
+
+        bfs_queue.push_back((NTreeNodeLocator::root(), 0));
+
+        while let Some((node_loc, explicit_loc)) = bfs_queue.pop_front()
+        {
+            // Iterate on children, that must exist
+            for (i, child_loc) in node_loc.children().iter().enumerate()
+            {
+                let child_data = self.map.get(child_loc).unwrap();
+                explicit_structure[explicit_loc].colors[i] = child_data.color;
+
+                if !child_data.subdivided
+                {
+                    explicit_structure[explicit_loc].structure[i] =
+                        StructurePointer::new(false, true, 0);
+                    continue;
+                }
+
+                // Child is subdivided, it needs to be handled.
+                let spot = explicit_structure.len();
+                explicit_structure.push(new_node.clone());
+                bfs_queue.push_back((*child_loc, spot));
+                explicit_structure[explicit_loc].structure[i] =
+                    StructurePointer::new(true, true, spot as u32);
+            }
+        }
+
+        (root.color, explicit_structure)
+    }
+}
+
+fn base_trailing_zeroes<const N: usize>(mut number: usize) -> Option<u32>
+{
+    if number == 0
+    {
+        return None;
+    }
+    let mut trailing_zeroes = 0;
+    while number % N == 0
+    {
+        number /= N;
+        trailing_zeroes += 1;
+    }
+
+    Some(trailing_zeroes)
+}
+
+impl<const N: usize> NTreeNodeLocator<N>
+{
+    pub fn root() -> Self
+    {
+        Self(1)
+    }
+
+    pub fn from_usize(n: usize) -> Self
+    {
+        Self(n)
+    }
+
+    pub fn new(mut depth: usize, x: usize, y: usize, z: usize) -> Self
+    {
+        let mut loc = Self::root();
+        let mut size = N.pow(depth as u32);
+        while depth != 0
+        {
+            size /= N;
+            let child_x = (x / size) % N;
+            let child_y = (y / size) % N;
+            let child_z = (z / size) % N;
+            loc = loc.child(child_x, child_y, child_z);
+            depth -= 1;
+        }
+
+        return loc;
+    }
+
+    pub fn depth(&self) -> usize
+    {
+        let mut n = self.0;
+        let mut depth = 0;
+        while n != 1
+        {
+            n /= N * N * N;
+            depth += 1;
+        }
+        depth
+    }
+
+    pub fn child_location(&self) -> (usize, usize, usize)
+    {
+        let mut n = self.0;
+        let z = n % N;
+        n /= N;
+        let y = n % N;
+        n /= N;
+        let x = n % N;
+
+        (x, y, z)
+    }
+
+    pub fn node_location(&self) -> (usize, usize, usize)
+    {
+        let mut locator = *self;
+        let mut nx = 0;
+        let mut ny = 0;
+        let mut nz = 0;
+        let mut size = 1;
+        while locator.0 != 1
+        {
+            let (cx, cy, cz) = locator.child_location();
+
+            nx += cx * size;
+            ny += cy * size;
+            nz += cz * size;
+            size *= N;
+
+            locator = locator.parent().unwrap();
+        }
+
+        (nx, ny, nz)
+    }
+
+    pub fn as_usize(&self) -> usize
+    {
+        self.0
+    }
+
+    pub fn parent(&self) -> Option<Self>
+    {
+        // Shift left by 3 digits in base N
+        let loc = self.0 / (N * N * N);
+        if loc == 0 { None } else { Some(Self(loc)) }
+    }
+
+    pub fn child(&self, x: usize, y: usize, z: usize) -> Self
+    {
+        assert!(x < N);
+        assert!(y < N);
+        assert!(z < N);
+
+        let child_loc = (self.0 * N * N * N) + (x * N * N) + (y * N) + z;
+        NTreeNodeLocator(child_loc)
+    }
+
+    pub fn children(&self) -> Vec<Self>
+    {
+        let mut children = vec![];
+
+        for x in 0..N
+        {
+            for y in 0..N
+            {
+                for z in 0..N
+                {
+                    children.push(self.child(x, y, z));
+                }
+            }
+        }
+        children
+    }
+}