diff --git a/Cargo.toml b/Cargo.toml
index d0c98c6..382199a 100644
--- a/Cargo.toml
+++ b/Cargo.toml
@@ -11,9 +11,13 @@ egui-wgpu = "0.36.1"
 egui-winit = "0.36.1"
 egui_plot = "0.37.0"
 env_logger = "0.11.11"
+fastapprox = "0.3.1"
 glam = "0.33.5"
+image = "0.25.10"
 itertools = "0.15.0"
 pollster = "1.0.1"
 rand = "0.10.2"
+rayon = "1.12.0"
+tiff = "0.11.3"
 wgpu = "30"
 winit = "0.30.13"
diff --git a/shaders/voxel.wgsl b/shaders/voxel.wgsl
index 2c131a6..68c2a3b 100644
--- a/shaders/voxel.wgsl
+++ b/shaders/voxel.wgsl
@@ -5,13 +5,15 @@ struct VertexOutput
     @location(1) color: vec4<f32>,
     @location(2) cam_pos: vec3<f32>,
     @location(3) world_pos: vec3<f32>,
+    @location(4) @interpolate(flat) structure_id: u32,
+    @location(5) chunk_position: vec3<f32>
 }
 
 struct ChunkImmediate
 {
     view_proj: mat4x4<f32>,
     cam_pos: vec3<f32>,
-    frame_timestamp: u32
+    frame_timestamp: u32,
 }
 
 var<immediate> constants: ChunkImmediate;
@@ -38,7 +40,7 @@ struct RequestBufferElement
 
 struct ColorPoolElement
 {
-    colors: array<vec4<f32>, 64>
+    colors: array<u32, 64>
 }
 
 struct LocationPoolElement
@@ -53,6 +55,16 @@ struct SortedRequestsElement
     child: u32
 }
 
+fn unpack_color(color: u32) -> vec4<f32>
+{
+    return vec4<f32>(
+            f32(color & 0xFF) / 255.,
+            f32((color >> 8) & 0xFF) / 255.,
+            f32((color >> 16) & 0xFF) / 255.,
+        f32((color >> 24) & 0xFF) / 255.
+    );
+}
+
 @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>;
@@ -61,8 +73,13 @@ struct SortedRequestsElement
 @group(0) @binding(5) var<storage, read_write> structure_table_pointer: array<u32>;
 @group(0) @binding(6) var<storage, read_write> structure_table_request_buffer: array<atomic<u32>>;
 
+struct FragmentOutput {
+    @location(0) color: vec4<f32>,
+    @builtin(frag_depth) depth: f32, // Equivalent to gl_FragDepth
+}
+
 @vertex
-fn chunk(@builtin(vertex_index) index: u32) -> VertexOutput
+fn chunk(@builtin(vertex_index) index: u32, @location(0) position: vec3<f32>, @location(1) id: u32) -> VertexOutput
 {
     let cube_vertices = array<vec3<f32>, 8>(
         vec3<f32>(0., 0., 0.),
@@ -98,14 +115,16 @@ fn chunk(@builtin(vertex_index) index: u32) -> VertexOutput
 
 
     let vertex = cube_vertices[cube_faces[quad_index * 4 + triangle_map[triangle_index]]];
-    let output_vertex = constants.view_proj * vec4<f32>(vertex, 1.0f);
+    let output_vertex = constants.view_proj * vec4<f32>(vertex + position, 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;
+    output.world_pos = vertex + position;
+    output.structure_id = id;
+    output.chunk_position = position;
 
     //let output = vec4<f32>(vertex, 1.0f);
     return output;
@@ -166,7 +185,7 @@ fn min_mask(x: vec3<f32>) -> vec3<bool>
 
 fn node_subdivided(node: u32) -> bool
 {
-    return (node >> 31) != 0;
+    return ((node >> 31) & 1) != 0;
 }
 
 fn node_pointer_valid(node: u32) -> bool
@@ -187,59 +206,228 @@ fn voxel_from_wall(position: vec3<f32>, ray_dir: vec3<f32>) -> vec3<i32>
     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>
+struct HitResult
 {
+    color: vec4<f32>,
+    hit_pos: vec3<f32>
+}
+
+fn new_traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_id: u32, dist_offset: f32) -> HitResult
+{
+    let max_depth = 5;
+    let dist_offset_voxel = dist_offset * f32(1 << u32(max_depth * 2));
+    let fovy_deg = 100.;
+    let cone_factor = tan((fovy_deg / 180.) * 3.14159) * 2.;
+
+    let st_pointer = structure_table_pointer[root_id];
+
+    
+    if (!node_subdivided(st_pointer))
+    {
+        var result: HitResult;
+        result.color = vec4(0., 1., 0., 1.);
+        result.hit_pos = ray_origin;
+        return result;
+    }
+    if(!node_pointer_valid(st_pointer))
+    {
+        // Node is subdivided, but not valid
+        // Send request on structure table
+        atomicAdd(&structure_table_request_buffer[root_id], 1); 
+
+        var result: HitResult;
+        result.color = vec4(0., 1., 0., 1.);
+        result.hit_pos = ray_origin;
+        return result;
+    }
+    //var current_node = node_pointer(st_pointer);
+
+    // Record usage
+    var dfs_stack = array<u32, 6>(node_pointer(st_pointer), 0, 0, 0, 0, 0);
+    var current_depth = 0;
+
+    usage_buffer[dfs_stack[current_depth]] = constants.frame_timestamp;
+
+    // Start location
+    //let voxel_dir = select(vec3(-1), vec3(1), ray_dir >= vec3(0.));
+    var node_size = 1 << u32(((max_depth - current_depth) * 2));
+    var child_size = node_size / 4;
+    var pos_origin = clamp(ray_origin * f32(1 << u32(max_depth * 2)), vec3(0.), vec3(f32(node_size) - 1.));
+    var voxel = vec3<i32>(pos_origin);
+    var far_t = 0.;
+
+    for(var iter = 0; iter < 400; iter ++)
+    {
+        // Compute child position
+        node_size = 1 << u32(((max_depth - current_depth) * 2));
+        child_size = node_size / 4;
+        var child_pos = (voxel / child_size) % 4;
+        var pointer = structure_pool[dfs_stack[current_depth]].pointers[child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4];
+
+        while(node_subdivided(pointer) && 
+            !((length(vec3<f32>(voxel) - pos_origin) + dist_offset) * cone_factor >= f32(node_size / 4))
+            )
+        {
+
+            if(!node_pointer_valid(pointer) && node_subdivided(pointer))
+            {
+                // Record request
+                atomicAdd(&request_buffer[dfs_stack[current_depth]].requests[child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4], 1);
+                break;
+            }
+
+            // Descend 
+            current_depth += 1;
+
+            node_size /= 4;
+            child_size /= 4;
+            child_pos = (voxel / child_size) % 4;
+            dfs_stack[current_depth] = node_pointer(pointer);
+
+            pointer = structure_pool[dfs_stack[current_depth]].pointers[child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4];
+
+            // Record usage
+            usage_buffer[dfs_stack[current_depth]] = constants.frame_timestamp;
+        }
+
+
+        // Check color
+        let color = color_pool[dfs_stack[current_depth]].colors[child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4];
+        if(((color >> 24) & 0xFF) != 0)
+        {
+            var result: HitResult;
+            result.color = unpack_color(color);
+            result.hit_pos = (far_t / f32(1 << u32(max_depth * 2))) * ray_dir + ray_origin;
+            return result;
+        }
+
+        // Advance
+        child_pos = (voxel / child_size) * child_size;
+        let far_wall = child_pos + select(vec3(0), vec3(child_size), ray_dir > vec3(0.));
+        let far_wall_inter = (vec3<f32>(far_wall) - pos_origin) / ray_dir;
+        far_t = min(min(far_wall_inter.x, far_wall_inter.y), far_wall_inter.z);
+
+        // Perform dda step on the children scale
+        let next_child = select(child_pos, child_pos + select(vec3(-1), vec3(1), ray_dir > vec3(0.)) * vec3(child_size), vec3(far_t) == far_wall_inter);
+        
+        let previous_voxel = voxel;
+        voxel = clamp(vec3<i32>(pos_origin + far_t * ray_dir), next_child, next_child + vec3(child_size) - vec3(1));
+
+        if any(voxel < vec3(0)) || any(voxel >= vec3(1 << u32((max_depth * 2))))
+        {
+            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 common_depth = ((countLeadingZeros(bit_diffs_lowest) - i32(32 - max_depth * 2)) / 2);
+
+        current_depth = common_depth;
+        //current_node = dfs_stack[current_depth];
+    }
+
+    // Iter max color
+    var result: HitResult;
+    result.color = vec4(1., 0., 1., 1.);
+    result.hit_pos = (far_t / f32(1 << u32(max_depth * 2))) * ray_dir + ray_origin;
+    return result;
+}
+
+fn traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_id: u32, dist_offset: f32) -> vec4<f32>
+{
+    let st_pointer = structure_table_pointer[root_id];
+
+    if (!node_subdivided(st_pointer))
+    {
+        return vec4(0., 1., 0., 1.);
+    }
+    if(!node_pointer_valid(st_pointer))
+    {
+        atomicAdd(&structure_table_request_buffer[root_id], 1); 
+        return vec4(0., 1., 0., 1.);
+    }
+
+    let fovy_deg = 100.;
+    let fovy = 3.14159 * (fovy_deg / 180.);
+    let definition = 1920.;
+    let cone_fovy = fovy / definition; 
+
+    let factor = 1.;
+    let cone_size_factor = 2. * tan(cone_fovy) * factor;
+
+    // Current depth of the node we are exploring
+    var current_depth = 0;
+
+    // Index of the current node's data
+    var current_node = u32(st_pointer & 0x3FFFFFFF);
+    usage_buffer[current_node] = constants.frame_timestamp;
+    var dfs_stack = array<u32, 6>(current_node, 0, 0, 0, 0, 0);
+
+
+
+    // Lut of the node_size per depth
+    var node_size_lut = array<i32, 6>(
+            4 * 4 * 4 * 4 * 4,
+            4 * 4 * 4 * 4,
+            4 * 4 * 4,
+            4 * 4,
+            4,
+            1,
+            );
+
+    let local_dist_offset = dist_offset * f32(node_size_lut[0]);
+
+
+    // Current node size
+    var node_size = node_size_lut[0]; // 128
+
+    // Size of a child of this node
+    var child_size = node_size / 4;
+
     // Simple FVT
     let t_off = abs(1. / ray_dir);
 
     // Start location
     let voxel_dir = select(vec3(-1), vec3(1), ray_dir >= vec3(0.));
-    var pos_origin = clamp(ray_origin * 256., vec3(0.), vec3(256. - 1.));
+    var pos_origin = clamp(ray_origin * f32(node_size), vec3(0.), vec3(f32(node_size) - 1.));
     var voxel = vec3<i32>(pos_origin);
     var last_voxel = voxel;
 
     let wall_offset = select(vec3(0), vec3(1), ray_dir > vec3(0.));
 
-    var dfs_stack = array<u32, 5>(0, 0, 0, 0, 0);
+    let max_depth = u32(5);
+    var adaptive_depth = i32(max_depth);
+    var far_t = 0.;
 
-        // Current depth of the node we are exploring
-        var current_depth = 0;
+    let ray_dir_inv = 1. / ray_dir;
+    let fma_offset = - pos_origin * ray_dir_inv;
 
-        // 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 ++)
+    //let depth_limit = 3;
+    for(var iter = 0; iter < 400; iter ++)
     {
 
         // Our ray is currently touching a voxel.
         // Descend to the lowest node that contains this voxel
 
         // Position of the child we are in
-        var child_pos = (vec3<u32>(voxel) >> vec3<u32>((4 - u32(current_depth + 1)) * 2)) & vec3<u32>(3); // Hardcode for 4-tree
+        var child_pos = (vec3<u32>(voxel) >> vec3<u32>((max_depth - u32(current_depth + 1)) * 2)) & vec3<u32>(3); // Hardcode for 4-tree
         var child_index = child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4;
 
         // Current node has been used. report
         usage_buffer[current_node] = constants.frame_timestamp;
         
         while(
-            node_subdivided(structure_pool[current_node].pointers[child_index]) 
-            && current_depth < depth_limit)
+            node_subdivided(structure_pool[current_node].pointers[child_index]) &&
+            (local_dist_offset + far_t) * cone_size_factor < f32(node_size_lut[current_depth])
+            )
         {
             if(!node_pointer_valid(structure_pool[current_node].pointers[child_index]))
             {
@@ -253,7 +441,7 @@ fn traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_color: vec4<f32>, ro
             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_pos = (vec3<u32>(voxel) >> vec3<u32>((max_depth - u32(current_depth + 1)) * 2)) & vec3<u32>(3); // Hardcode for 4-tree
             child_index = child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4;
             child_size = node_size / 4;
         }
@@ -267,23 +455,26 @@ fn traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_color: vec4<f32>, ro
         // It is guaranteed that the child is leave
 
         // Check current leave's color
-        let color = color_pool[current_node].colors[child_index];
+        let color = unpack_color(color_pool[current_node].colors[child_index]);
         if(color.w != 0.) // Not transparent
         {
+            /*
             let k = child_pos.x + child_pos.y + child_pos.z;
             let w = voxel.x + voxel.y + voxel.z;
             let x = select(0.5, 1., k % 2 == 0) * select(0.8, 1., w % 2 == 0);
 
             var div = 1;
             var overlay = 1.;
-            for(var i = 1; i <= 4; i++)
+            for(var i = 1; i <= 5; i++)
             {
                 let x = (voxel.x / div + voxel.y / div + voxel.z / div) % 2 == 0;
                 overlay -= select(0., 1. / (f32(i) * 2.5), x);
                 div *= 4;
             }
+            */
 
-            return overlay * color;
+            return color;
+            //return overlay * color;
         }
 
         // Voxel and whole child containing it is empty
@@ -291,16 +482,17 @@ fn traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_color: vec4<f32>, ro
         // 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 far_ts = (vec3<f32>(far_corner) - pos_origin) / ray_dir; // TODO: Turn into fma
+        let far_ts = fma(vec3<f32>(far_corner), ray_dir_inv, fma_offset);
+        far_t = min(min(far_ts.x, far_ts.y), far_ts.z);
 
-        let next_child_min = select(child_position, child_position + wall_offset * child_size, vec3(far_t) == far_ts);
-        let next_child_max = next_child_min + vec3(child_size);
+        let next_child_min = select(child_position, child_position + voxel_dir * child_size, vec3(far_t) == far_ts);
+        let next_child_max = next_child_min + vec3(child_size) - vec3(1);
         
         // The ray (far_t) is now touching the new child to explore
         // Find out which actual voxel we are touching
         let previous_voxel = voxel;
-        let float_voxel = clamp(pos_origin + far_t * ray_dir, vec3<f32>(next_child_min), vec3<f32>(next_child_max));
+        let float_voxel = clamp(vec3<i32>(pos_origin + far_t * ray_dir), next_child_min, next_child_max);
         /*
         voxel = vec3<i32>(
             floor(
@@ -312,8 +504,10 @@ fn traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_color: vec4<f32>, ro
         );
         */
         //voxel = vec3<i32>(round(float_voxel));
-        voxel = voxel_from_wall(float_voxel, ray_dir);
-        if(any(voxel < vec3(0)) || any(voxel >= vec3(256)))
+        //voxel = voxel_from_wall(float_voxel, ray_dir);
+        //voxel = voxel_from_wall(float_voxel, ray_dir);
+        voxel = float_voxel;
+        if(any(voxel < vec3(0)) || any(voxel >= vec3(node_size_lut[0])))
         {
             //return vec4(f32(iter) / 100.);
             discard;
@@ -329,7 +523,7 @@ fn traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_color: vec4<f32>, ro
         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 flb = ((countLeadingZeros(bit_diffs_lowest) - i32(32 - max_depth * 2)) / 2);
         let common_depth = flb;
 
         current_depth = common_depth;
@@ -345,30 +539,27 @@ fn traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_color: vec4<f32>, ro
 
 }
 
+@early_depth_test(less_equal)
 @fragment
-fn fragment(in: VertexOutput) -> @location(0) vec4<f32>
+fn fragment(in: VertexOutput) -> FragmentOutput
 {
+    //frag_out.color = vec4<f32>(2 * 0.01 / (100. + 0.01 - depth * (100. - 0.01)));
     let ray_dir = normalize(in.world_pos - in.cam_pos);
-    let interp = box_inter(in.cam_pos, ray_dir, vec3(0.), vec3(1));
-    let ray_origin = in.cam_pos + ray_dir * (max(0., interp.x));
+    let interp = box_inter(in.cam_pos - in.chunk_position, ray_dir, vec3(0.), vec3(1));
+    let ray_origin = (in.cam_pos - in.chunk_position) + ray_dir * (max(0., interp.x));
 
-    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.);
+    let result = new_traverse(ray_dir, ray_origin, in.structure_id, length(in.cam_pos - ray_origin));
+    let clip_pos = constants.view_proj * vec4(result.hit_pos + in.chunk_position, 1.);
+    let depth = clip_pos.z / clip_pos.w;
+    var frag_out: FragmentOutput;
+    frag_out.color = result.color;
+    frag_out.depth = depth;
+    return frag_out;
+
+    //return vec4<f32>(ray_origin, 1.);
+    //return frag_out;
+    //return vec4(interp.y / 10.);
 }
 
 /*
diff --git a/src/camera.rs b/src/camera.rs
index 5dc122e..25a4ec9 100644
--- a/src/camera.rs
+++ b/src/camera.rs
@@ -18,7 +18,7 @@ pub struct Camera
     fov: f32,
     pub aspect: f32,
 
-    speed: f32,
+    pub speed: f32,
     pub pressed_keyset: HashSet<KeyCode>,
 }
 
diff --git a/src/main.rs b/src/main.rs
index ac2eb9b..50d76b3 100644
--- a/src/main.rs
+++ b/src/main.rs
@@ -4,6 +4,7 @@ use core::sync;
 use std::cell::RefCell;
 use std::collections::HashMap;
 use std::fs::File;
+use std::hash::Hash;
 use std::rc::Rc;
 use std::sync::Arc;
 use std::sync::atomic::AtomicBool;
@@ -13,6 +14,8 @@ use bytemuck::Pod;
 use bytemuck::Zeroable;
 use crevice::std140::AsStd140;
 use crevice::std430::AsStd430;
+use egui::Color32;
+use egui::Label;
 use egui::emath::fast_midpoint;
 use egui::mutex::Mutex;
 use egui_plot::BarChart;
@@ -21,6 +24,9 @@ use glam::Vec3;
 use glam::Vec4;
 use itertools::Itertools;
 use rand::random;
+use rayon::iter::IndexedParallelIterator;
+use rayon::iter::IntoParallelRefIterator;
+use rayon::iter::ParallelIterator;
 use wgpu::BindGroup;
 use wgpu::BindGroupEntry;
 use wgpu::BindGroupLayoutDescriptor;
@@ -54,6 +60,7 @@ use wgpu::util::DownloadBuffer;
 use wgpu::util::StagingBelt;
 use winit::application::ApplicationHandler;
 use winit::event::DeviceEvent;
+use winit::event::MouseScrollDelta;
 use winit::event::WindowEvent;
 use winit::event_loop;
 use winit::event_loop::ActiveEventLoop;
@@ -66,9 +73,14 @@ use winit::window::WindowId;
 
 use crate::camera::Camera;
 use crate::egui_renderer::EguiRenderer;
+use crate::producers::BallGenerator;
+use crate::producers::ChunkedProducer;
+use crate::producers::Producer;
 use crate::producers::SineGenerator;
+use crate::producers::TerrainGenerator;
 use crate::voxel::cache::CacheNodeRequest;
 use crate::voxel::cache::CacheResponse;
+use crate::voxel::cache::ColorBytes;
 use crate::voxel::cache::DestinationElement;
 use crate::voxel::cache::LocationPoolElement;
 use crate::voxel::cache::RequestBuffer;
@@ -103,6 +115,10 @@ struct State
 
     pipeline: RenderPipeline,
     voxel_cache: Arc<Mutex<VoxelCache<4>>>,
+    terrain_generator: Arc<TerrainGenerator<4>>,
+    chunk_pos_map: Arc<HashMap<u32, (usize, usize, usize)>>,
+    instance_buffer: Buffer,
+    instance_count: usize,
     usage_vec: Arc<Mutex<Vec<usize>>>,
     insertion_debounce: bool,
 
@@ -123,6 +139,16 @@ struct Immediates
     frame_timestamp: u32,
 }
 
+#[derive(Debug, Clone, Copy, Zeroable, Pod)]
+#[repr(C)]
+struct InstanceAttribute
+{
+    x: f32,
+    y: f32,
+    z: f32,
+    id: u32,
+}
+
 impl State
 {
     async fn new(display: OwnedDisplayHandle, window: Arc<Window>) -> State
@@ -141,7 +167,7 @@ impl State
             .unwrap();
         let (device, queue) = adapter
             .request_device(&wgpu::DeviceDescriptor {
-                required_features: Features::IMMEDIATES,
+                required_features: Features::IMMEDIATES | Features::SHADER_EARLY_DEPTH_TEST,
                 required_limits: wgpu::Limits {
                     max_immediate_size: 112,
                     max_storage_buffers_per_shader_stage: 16,
@@ -160,9 +186,34 @@ impl State
 
         let egui_renderer = EguiRenderer::new(&device, surface_format, &window);
 
-        let mut voxel_cache = VoxelCache::<4>::new(16000, device.clone(), queue.clone());
-        let id = voxel_cache.structure_table.allocate_structure(true);
-        println!("id: {id}");
+        let mut voxel_cache = VoxelCache::<4>::new(100_000, device.clone(), queue.clone());
+
+        let terrain_generator = TerrainGenerator::<4>::new(5, "vxls_height.tif", 0.2, "img.jpg");
+
+        let mut chunk_pos_map = HashMap::new();
+        let chunk_instances = (0..terrain_generator.chunk_width)
+            .cartesian_product(0..terrain_generator.chunk_height)
+            .cartesian_product(0..terrain_generator.chunk_alt)
+            .map(|((x, z), y)| {
+                let chunk_id = voxel_cache.structure_table.allocate_structure(true);
+                //dbg!(chunk_id);
+
+                chunk_pos_map.insert(chunk_id, (x, y, z));
+                InstanceAttribute {
+                    x: x as f32,
+                    y: y as f32,
+                    z: z as f32,
+                    id: chunk_id,
+                }
+            })
+            .collect::<Vec<_>>();
+        let instance_count = chunk_instances.len();
+
+        let instance_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
+            label: Some("Instance buffer"),
+            contents: bytemuck::cast_slice(chunk_instances.as_slice()),
+            usage: BufferUsages::COPY_DST | BufferUsages::VERTEX,
+        });
 
         let shader_module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
             label: Some("Main shader module"),
@@ -187,13 +238,28 @@ impl State
                 module: &shader_module,
                 entry_point: Some("chunk"),
                 compilation_options: Default::default(),
-                buffers: &[],
+                buffers: &[Some(wgpu::VertexBufferLayout {
+                    array_stride: (size_of::<f32>() * 3 + size_of::<u32>()) as u64,
+                    step_mode: wgpu::VertexStepMode::Instance,
+                    attributes: &[
+                        wgpu::VertexAttribute {
+                            format: wgpu::VertexFormat::Float32x3,
+                            offset: 0,
+                            shader_location: 0,
+                        },
+                        wgpu::VertexAttribute {
+                            format: wgpu::VertexFormat::Uint32,
+                            offset: (size_of::<f32>() * 3) as u64,
+                            shader_location: 1,
+                        },
+                    ],
+                })],
             },
             primitive: wgpu::PrimitiveState {
                 topology: wgpu::PrimitiveTopology::TriangleList,
                 strip_index_format: None,
                 front_face: wgpu::FrontFace::Ccw,
-                cull_mode: None,
+                cull_mode: Some(wgpu::Face::Front),
                 unclipped_depth: false,
                 polygon_mode: wgpu::PolygonMode::Fill,
                 conservative: false,
@@ -201,7 +267,7 @@ impl State
             depth_stencil: Some(wgpu::DepthStencilState {
                 format: wgpu::TextureFormat::Depth24PlusStencil8,
                 depth_write_enabled: Some(true),
-                depth_compare: Some(wgpu::CompareFunction::LessEqual),
+                depth_compare: Some(wgpu::CompareFunction::Less),
                 stencil: wgpu::StencilState::default(),
                 bias: wgpu::DepthBiasState::default(),
             }),
@@ -235,6 +301,10 @@ impl State
             voxel_cache: Arc::new(Mutex::new(voxel_cache)),
             insertion_debounce: false,
             camera: Default::default(),
+            instance_buffer,
+            instance_count,
+            terrain_generator: Arc::new(terrain_generator),
+            chunk_pos_map: chunk_pos_map.into(),
         };
 
         // Configure surface for the first time
@@ -323,9 +393,7 @@ impl State
     fn render(&mut self)
     {
         self.camera.update();
-
-        // Write random shit in request buffer
-        let mut belt = StagingBelt::new(self.device.clone(), size_of::<u32>() as u64);
+        self.voxel_cache.lock().next_frame();
 
         // Create texture view.
         // NOTE: We must handle Timeout because the surface may be unavailable
@@ -395,6 +463,7 @@ impl State
             });
 
             renderpass.set_pipeline(&self.pipeline);
+            renderpass.set_vertex_buffer(0, self.instance_buffer.slice(..));
             renderpass.set_bind_group(0, Some(&self.voxel_cache.lock().bind_group()), &[]);
             let imm = [Immediates {
                 view_proj: self.camera.view_proj(),
@@ -402,7 +471,7 @@ impl State
                 frame_timestamp: self.voxel_cache.lock().current_timestamp(),
             }];
             renderpass.set_immediates(0, unsafe { as_raw_bytes(&imm) });
-            renderpass.draw(0..36, 0..1);
+            renderpass.draw(0..36, 0..(self.instance_count as u32));
 
             // End the renderpass.
             drop(renderpass);
@@ -410,14 +479,20 @@ impl State
 
         let requests = self.device.create_buffer(&wgpu::BufferDescriptor {
             label: Some("dummy_dumb_dinky_aaaahhh_buffer"),
-            size: 16 * 256,
+            size: 16 * 1024,
             usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
             mapped_at_creation: false,
         });
 
-        self.voxel_cache
-            .lock()
-            .cache_post_render(&mut encoder, requests.clone());
+        if !self
+            .camera
+            .pressed_keyset
+            .contains(&winit::keyboard::KeyCode::KeyF)
+        {
+            self.voxel_cache
+                .lock()
+                .cache_post_render(&mut encoder, requests.clone());
+        }
 
         // If you wanted to call any drawing commands, they would go here.
         {
@@ -425,6 +500,17 @@ impl State
             egui::Window::new("Window ! ").resizable(true).show(
                 self.egui_renderer.context(),
                 |ui| {
+                    if self
+                        .camera
+                        .pressed_keyset
+                        .contains(&winit::keyboard::KeyCode::KeyF)
+                    {
+                        ui.label(
+                            egui::RichText::new("Cache paused")
+                                .color(Color32::RED)
+                                .size(28.),
+                        );
+                    }
                     egui_plot::Plot::new("Plot").show(ui, |plot_ui| {
                         plot_ui.bar_chart(BarChart::new(
                             "histo",
@@ -488,17 +574,24 @@ impl State
             self.insertion_debounce = false;
         }
 
-        if (self
+        // if (self
+        //     .camera
+        //     .pressed_keyset
+        //     .contains(&winit::keyboard::KeyCode::KeyF))
+        //     && !self.insertion_debounce
+        // {
+        if !self
             .camera
             .pressed_keyset
-            .contains(&winit::keyboard::KeyCode::KeyF))
-            && !self.insertion_debounce
+            .contains(&winit::keyboard::KeyCode::KeyF)
         {
             self.insertion_debounce = true;
             let request_count = self.voxel_cache.lock().total_request_count();
             let cloned_cache = self.voxel_cache.clone();
             let cloned_device = self.device.clone();
             let cloned_queue = self.queue.clone();
+            let cloned_generator = self.terrain_generator.clone();
+            let cloned_map = self.chunk_pos_map.clone();
             let (tx, rx) = sync_channel(1);
             DownloadBuffer::read_buffer(
                 &self.device.clone(),
@@ -513,58 +606,88 @@ impl State
                     let cache_node_requests: Vec<CacheNodeRequest> =
                         bytemuck::pod_collect_to_vec(&buffer.unwrap());
 
-                    let mut sine_gen = SineGenerator::<4>::new(4);
+                    let generator = cloned_generator;
+                    let gen_test = SineGenerator::<4>::new(5);
+
                     let mut structure_nodes = vec![];
-                    let mut color_nodes = vec![];
+                    let mut color_nodes: Vec<[ColorBytes; 64]> = vec![];
                     let mut location_nodes = vec![];
                     let mut destinations = vec![];
-                    for request in cache_node_requests.iter().take(request_count as usize)
-                    {
-                        let location;
-                        let node;
-                        if request.child_index == u32::MAX
-                        {
-                            // Produce root node
-                            node = sine_gen.produce_node(0, 0, 0, 0);
-                            location = LocationPoolElement {
-                                structure_id: 0,
-                                structure_locator: NTreeNodeLocator::<4>::root().as_usize() as u32,
-                            };
-                        }
-                        else
-                        {
-                            // Figure out depth of request
-                            let locator = NTreeNodeLocator::<4>::from_usize(
-                                request.structure_locator as usize,
-                            );
-                            let depth = locator.depth();
-                            //let (x, y, z) = locator.node_location();
 
-                            let child_x = request.child_index / (4 * 4);
-                            let child_y = (request.child_index % (4 * 4)) / 4;
-                            let child_z = request.child_index % 4;
+                    cache_node_requests
+                        .par_iter()
+                        .take(request_count as usize)
+                        .map(|request| {
+                            // Unpack structure id
+                            let (chunk_x, chunk_y, chunk_z) =
+                                *cloned_map.get(&request.structure_id).unwrap();
 
-                            let child_locator =
-                                locator.child(child_x as usize, child_y as usize, child_z as usize);
+                            let location;
+                            let node;
+                            if request.child_index == u32::MAX
+                            {
+                                // Produce root node
+                                node =
+                                    generator.produce_node(0, 0, 0, 0, (chunk_x, chunk_y, chunk_z));
+                                location = LocationPoolElement {
+                                    structure_id: request.structure_id,
+                                    structure_locator: NTreeNodeLocator::<4>::root().as_usize()
+                                        as u32,
+                                };
+                            }
+                            else
+                            {
+                                // Figure out depth of request
+                                let locator = NTreeNodeLocator::<4>::from_usize(
+                                    request.structure_locator as usize,
+                                );
+                                let depth = locator.depth();
+                                //let (x, y, z) = locator.node_location();
 
-                            let (nx, ny, nz) = child_locator.node_location();
+                                let child_z = request.child_index / (4 * 4);
+                                let child_y = (request.child_index % (4 * 4)) / 4;
+                                let child_x = request.child_index % 4;
 
-                            node = sine_gen.produce_node(depth + 1, nx, ny, nz);
+                                let child_locator = locator.child(
+                                    child_x as usize,
+                                    child_y as usize,
+                                    child_z as usize,
+                                );
 
-                            location = LocationPoolElement {
-                                structure_id: 0,
-                                structure_locator: child_locator.as_usize() as u32,
-                            };
-                        }
+                                let (nx, ny, nz) = child_locator.node_location();
 
-                        structure_nodes.push(node.structure);
-                        color_nodes.push(node.colors);
-                        location_nodes.push(location);
-                        destinations.push(DestinationElement {
-                            node: request.node_index,
-                            child: request.child_index,
+                                node = generator.produce_node(
+                                    depth + 1,
+                                    nx,
+                                    ny,
+                                    nz,
+                                    (chunk_x, chunk_y, chunk_z),
+                                );
+
+                                location = LocationPoolElement {
+                                    structure_id: request.structure_id,
+                                    structure_locator: child_locator.as_usize() as u32,
+                                };
+                            }
+
+                            (
+                                node.structure,
+                                node.colors,
+                                location,
+                                DestinationElement {
+                                    node: request.node_index,
+                                    child: request.child_index,
+                                },
+                            )
+                        })
+                        .collect::<Vec<_>>()
+                        .into_iter()
+                        .for_each(|data| {
+                            structure_nodes.push(data.0);
+                            color_nodes.push(std::array::from_fn(|i| data.1[i].into()));
+                            location_nodes.push(data.2);
+                            destinations.push(data.3);
                         });
-                    }
 
                     if structure_nodes.len() != 0
                     {
@@ -607,7 +730,7 @@ impl State
                     }
                 },
             );
-            println!("Total request count: {}", request_count);
+            //println!("Total request count: {}", request_count);
             loop
             {
                 if rx.try_recv().is_ok()
@@ -617,10 +740,17 @@ impl State
                 let _ = self.device.poll(wgpu::wgt::PollType::Poll);
             }
 
-            self.voxel_cache.lock().next_frame();
             drop(rx);
         }
     }
+
+    pub fn mouse_wheel(&mut self, delta: MouseScrollDelta)
+    {
+        if let MouseScrollDelta::LineDelta(_, y) = delta
+        {
+            self.camera.speed += y * (self.camera.speed * 0.05);
+        }
+    }
 }
 
 #[derive(Default)]
@@ -679,6 +809,11 @@ impl ApplicationHandler for App
                 // here as this event is always followed up by redraw request.
                 state.resize(size);
             }
+            WindowEvent::MouseWheel { delta, .. } =>
+            {
+                state.mouse_wheel(delta);
+            }
+
             _ => (),
         }
     }
diff --git a/src/producers.rs b/src/producers.rs
index e6b5386..209c85b 100644
--- a/src/producers.rs
+++ b/src/producers.rs
@@ -1,3 +1,6 @@
+use std::fs::File;
+use std::path::Path;
+
 use glam::Vec3;
 use itertools::Itertools;
 
@@ -5,7 +8,7 @@ use crate::voxel::gpu::ExplicitNTreeNode;
 use crate::voxel::gpu::StructurePointer;
 use crate::voxel::sparse::Color;
 
-pub struct SineGenerator<const N: usize>
+pub struct BallGenerator<const N: usize>
 {
     chunk_power: usize,
 }
@@ -15,22 +18,28 @@ pub fn map(x: f32, x_min: f32, x_max: f32, y_min: f32, y_max: f32) -> f32
     ((x - x_min) / (x_max - x_min)) * (y_max - y_min) + y_min
 }
 
-impl<const N: usize> SineGenerator<N>
+pub trait Producer<const N: usize>
+where
+    [(); N * N * N]:,
+{
+    fn produce_node(&self, depth: usize, nx: usize, ny: usize, nz: usize) -> ExplicitNTreeNode<N>;
+}
+
+impl<const N: usize> BallGenerator<N>
 where
     [(); N * N * N]:,
 {
     pub fn new(chunk_power: usize) -> Self
     {
-        SineGenerator { chunk_power }
+        BallGenerator { chunk_power }
     }
+}
 
-    pub fn produce_node(
-        &self,
-        depth: usize,
-        nx: usize,
-        ny: usize,
-        nz: usize,
-    ) -> ExplicitNTreeNode<N>
+impl<const N: usize> Producer<N> for BallGenerator<N>
+where
+    [(); N * N * N]:,
+{
+    fn produce_node(&self, depth: usize, nx: usize, ny: usize, nz: usize) -> ExplicitNTreeNode<N>
     {
         let node_size = N.pow((self.chunk_power - depth) as u32);
         let child_size = node_size / N;
@@ -49,23 +58,38 @@ where
             let gvy = gny + cy * child_size + (child_size / 2);
             let gvz = gnz + cz * child_size + (child_size / 2);
 
-            let dist = Vec3::new(gvx as f32 - 128., gvy as f32 - 128., gvz as f32 - 128.).length();
+            let dist = Vec3::new(
+                gvx as f32 - (global_size / 2) as f32,
+                gvy as f32 - (global_size / 2) as f32,
+                gvz as f32 - (global_size / 2) as f32,
+            )
+            .length();
             let child_diagonal_length = (child_size as f32 / 2.) * f32::sqrt(3.);
 
             let alpha;
-            if (dist - 128.).abs() <= child_diagonal_length
+            if (dist - (global_size as f32 / 2.)).abs() <= child_diagonal_length
             {
-                children.push(StructurePointer::new(depth <= 3, false, 0));
+                //children.push(StructurePointer::new(depth <= 3, false, 0));
+                children.push(StructurePointer(
+                    if depth < (self.chunk_power - 1)
+                    {
+                        0xFFFFFFFF
+                    }
+                    else
+                    {
+                        0
+                    },
+                ));
                 alpha = if dist > 128. { 0. } else { 1. };
             }
-            else if dist > 128.
+            else if dist > (global_size as f32 / 2.)
             {
-                children.push(StructurePointer::new(false, false, 0));
+                children.push(StructurePointer(0));
                 alpha = 0.;
             }
             else
             {
-                children.push(StructurePointer::new(false, false, 0));
+                children.push(StructurePointer(0));
                 alpha = 1.;
             }
 
@@ -86,3 +110,340 @@ where
         }
     }
 }
+
+pub struct SineGenerator<const N: usize>
+{
+    chunk_power: usize,
+}
+
+impl<const N: usize> SineGenerator<N>
+where
+    [(); N * N * N]:,
+{
+    pub fn new(chunk_power: usize) -> Self
+    {
+        Self { chunk_power }
+    }
+}
+
+impl<const N: usize> Producer<N> for SineGenerator<N>
+where
+    [(); N * N * N]:,
+{
+    fn produce_node(&self, depth: usize, nx: usize, ny: usize, nz: usize) -> ExplicitNTreeNode<N>
+    {
+        let node_size = N.pow((self.chunk_power - depth) as u32);
+        let child_size = node_size / N;
+        let global_size = N.pow(self.chunk_power as u32);
+
+        let mut children = vec![];
+        let mut children_color = vec![];
+
+        let gnx = nx * node_size;
+        let gny = ny * node_size;
+        let gnz = nz * node_size;
+        // Iterate on children of this node
+        for ((cx, cy), cz) in (0..N).cartesian_product(0..N).cartesian_product(0..N)
+        {
+            let total_sub_voxels = child_size * child_size * child_size;
+            let mut filled_sub_voxels = 0;
+
+            let gcx = gnx + cx * child_size;
+            let gcy = gny + cy * child_size;
+            let gcz = gnz + cz * child_size;
+
+            // prepare 2d values
+
+            // Iterate on children
+            for (x, z) in (0..child_size).cartesian_product(0..child_size)
+            {
+                let gvx = gcx + x;
+                let gvz = gcz + z;
+                //let gvy = gcy;
+                let sx = map(gvx as f32, 0., global_size as f32, -8., 8.).abs();
+                let sz = map(gvz as f32, 0., global_size as f32, -8., 8.).abs();
+                let sample = (fastapprox::fast::cos(sx) + fastapprox::fast::cos(sz)) * 0.5;
+                let sample_height = map(sample, -1., 1., 0., 500.);
+
+                let prop = map(
+                    sample_height,
+                    gcy as f32,
+                    (gcy + child_size) as f32,
+                    0.,
+                    child_size as f32,
+                )
+                .clamp(0., child_size as f32)
+                .floor() as usize;
+                filled_sub_voxels += prop;
+            }
+
+            let alpha;
+            if filled_sub_voxels == 0
+            {
+                children.push(StructurePointer(0));
+                alpha = 0.;
+            }
+            else if filled_sub_voxels >= total_sub_voxels
+            {
+                children.push(StructurePointer(0));
+                alpha = 1.;
+            }
+            else
+            {
+                children.push(StructurePointer(
+                    if depth < (self.chunk_power - 1)
+                    {
+                        0xFFFFFFFF
+                    }
+                    else
+                    {
+                        0
+                    },
+                ));
+                //children.push(StructurePointer(0));
+                alpha = if filled_sub_voxels > total_sub_voxels / 2
+                {
+                    1.
+                }
+                else
+                {
+                    0.
+                };
+            }
+
+            children_color.push(Color(
+                (gcx + child_size / 2) as f32 / global_size as f32,
+                (gcy + child_size / 2) as f32 / global_size as f32,
+                (gcz + child_size / 2) as f32 / global_size as f32,
+                alpha,
+            ));
+        }
+
+        ExplicitNTreeNode {
+            structure: std::array::from_fn(|i| children[i]),
+            colors: std::array::from_fn(|i| children_color[i]),
+        }
+    }
+}
+
+pub struct TerrainGenerator<const N: usize>
+where
+    [(); N * N * N]:,
+{
+    chunk_power: usize,
+    heightmap_width: usize,
+    heightmap_height: usize,
+    terrain_width: usize,
+    terrain_height: usize,
+    heightmap_min: f32,
+    heightmap_max: f32,
+    heightmap: Vec<f32>,
+    colormap: Vec<u8>,
+
+    pub chunk_width: usize,
+    pub chunk_height: usize,
+    pub chunk_alt: usize,
+}
+
+impl<const N: usize> TerrainGenerator<N>
+where
+    [(); N * N * N]:,
+{
+    pub fn new<P: AsRef<Path>>(
+        chunk_power: usize,
+        height_path: P,
+        height_factor: f32,
+        color_path: P,
+    ) -> Self
+    {
+        let mut tiff_dec = tiff::decoder::Decoder::new(File::open(height_path).unwrap()).unwrap();
+        let (heightmap_width, heightmap_height) = tiff_dec.dimensions().unwrap();
+        let (heightmap_width, heightmap_height) =
+            (heightmap_width as usize, heightmap_height as usize);
+
+        let heightmap = match tiff_dec.read_image().unwrap()
+        {
+            tiff::decoder::DecodingResult::F32(vec) => vec,
+            _ => panic!("Unsupported format"),
+        };
+
+        let mut color = image::ImageReader::open(color_path).unwrap();
+        color.no_limits();
+
+        let color = color.decode().unwrap();
+        let colormap = color.as_rgb8().unwrap().to_vec();
+
+        let terrain_width = color.width() as usize;
+        let terrain_height = color.height() as usize;
+
+        let heightmap_min = heightmap.iter().copied().reduce(f32::min).unwrap();
+        let heightmap_max = heightmap.iter().copied().reduce(f32::max).unwrap();
+
+        // Decide size in chunks
+        let height_amplitude = heightmap_max - heightmap_min;
+        let chunk_size = N.pow(chunk_power as u32);
+        let chunk_width = terrain_width.div_ceil(chunk_size);
+        let chunk_height = terrain_height.div_ceil(chunk_size);
+        let chunk_alt = ((height_amplitude / height_factor) as usize).div_ceil(chunk_size);
+
+        Self {
+            chunk_power,
+            heightmap_width,
+            heightmap_height,
+            heightmap_min,
+            heightmap_max,
+            terrain_width,
+            terrain_height,
+            heightmap,
+            colormap,
+
+            chunk_width,
+            chunk_height,
+            chunk_alt,
+        }
+    }
+}
+
+impl<const N: usize> ChunkedProducer<N> for TerrainGenerator<N>
+where
+    [(); N * N * N]:,
+{
+    fn produce_node(
+        &self,
+        depth: usize,
+        nx: usize,
+        ny: usize,
+        nz: usize,
+        chunk_pos: (usize, usize, usize),
+    ) -> ExplicitNTreeNode<N>
+    {
+        let node_size = N.pow((self.chunk_power - depth) as u32);
+        let child_size = node_size / N;
+        let global_size = N.pow(self.chunk_power as u32);
+
+        let mut children = vec![StructurePointer(0); N * N * N];
+        let mut children_color = vec![Color(0., 0., 0., 0.); N * N * N];
+
+        let gnx = chunk_pos.0 * global_size + nx * node_size;
+        let gny = chunk_pos.1 * global_size + ny * node_size;
+        let gnz = chunk_pos.2 * global_size + nz * node_size;
+        // Iterate on children of this node
+        for (cx, cz) in (0..N).cartesian_product(0..N)
+        {
+            let gcx = gnx + cx * child_size;
+            let gcz = gnz + cz * child_size;
+
+            // prepare 2d values
+
+            // Iterate on children
+            let mut sample_max = self.heightmap_min;
+            let mut sample_min = self.heightmap_max;
+            let mut color_avg = Color(0., 0., 0., 0.);
+            let mut count = 0;
+            for (x, z) in (0..child_size).cartesian_product(0..child_size)
+            {
+                let gvx = gcx + x;
+                let gvz = gcz + z;
+                if gvx < self.terrain_width && gvz < self.terrain_height
+                {
+                    // Height sample
+                    let height_x = (gvx * self.heightmap_width) / self.terrain_width;
+                    let height_z = (gvz * self.heightmap_height) / self.terrain_height;
+
+                    let sample = self.heightmap[height_x + height_z * self.heightmap_width];
+                    sample_max = sample_max.max(sample);
+                    sample_min = sample_min.min(sample);
+
+                    let sample_color_r = self.colormap[(gvx + gvz * self.terrain_width) * 3];
+                    let sample_color_g = self.colormap[(gvx + gvz * self.terrain_width) * 3 + 1];
+                    let sample_color_b = self.colormap[(gvx + gvz * self.terrain_width) * 3 + 2];
+
+                    color_avg.0 += map(sample_color_r as f32, 0., 256., 0., 1.);
+                    color_avg.1 += map(sample_color_g as f32, 0., 256., 0., 1.);
+                    color_avg.2 += map(sample_color_b as f32, 0., 256., 0., 1.);
+                    count += 1;
+                }
+                //let gvy = gcy;
+            }
+
+            color_avg.0 /= count as f32;
+            color_avg.1 /= count as f32;
+            color_avg.2 /= count as f32;
+
+            let sample_min = map(
+                sample_min,
+                self.heightmap_min,
+                self.heightmap_max,
+                0.,
+                (self.chunk_alt * global_size) as f32,
+            ) as usize;
+            let sample_max = map(
+                sample_max,
+                self.heightmap_min,
+                self.heightmap_max,
+                0.,
+                (self.chunk_alt * global_size) as f32,
+            ) as usize;
+
+            for cy in 0..N
+            {
+                let gcy = gny + cy * child_size;
+                let index = cz * N * N + cy * N + cx;
+                let alpha;
+                if gcy > sample_max
+                {
+                    children[index] = StructurePointer(0);
+                    alpha = 0.;
+                }
+                else if gcy + child_size < sample_min
+                {
+                    children[index] = StructurePointer(0);
+                    alpha = 1.;
+                }
+                else
+                {
+                    children[index] = StructurePointer(
+                        if depth < (self.chunk_power - 1)
+                        {
+                            0xFFFFFFFF
+                        }
+                        else
+                        {
+                            0
+                        },
+                    );
+                    //children.push(StructurePointer(0));
+                    alpha = if (sample_max + sample_min / 2) > gcy + (child_size / 2)
+                    {
+                        1.
+                    }
+                    else
+                    {
+                        0.
+                    };
+                }
+
+                children_color[index] = Color(color_avg.0, color_avg.1, color_avg.2, alpha);
+            }
+        }
+
+        ExplicitNTreeNode {
+            structure: std::array::from_fn(|i| children[i]),
+            colors: std::array::from_fn(|i| children_color[i]),
+        }
+    }
+}
+
+pub trait ChunkedProducer<const N: usize>
+where
+    [(); N * N * N]:,
+{
+    fn produce_node(
+        &self,
+        depth: usize,
+        nx: usize,
+        ny: usize,
+        nz: usize,
+        chunk_pos: (usize, usize, usize),
+    ) -> ExplicitNTreeNode<N>;
+}
diff --git a/src/voxel/cache.rs b/src/voxel/cache.rs
index d96ea10..cf5ca35 100644
--- a/src/voxel/cache.rs
+++ b/src/voxel/cache.rs
@@ -209,32 +209,39 @@ where
 
 
                             @compute
-                            @workgroup_size(16)
+                            @workgroup_size(64)
                             fn main(
                                     @builtin(global_invocation_id) global_invocation_id: vec3<u32>
                                 )
                             {{
-                                 let request_count_round_up =
-                                     ((request_count / {children_count}) + select(u32(0), u32(1), request_count % {children_count} != 0)) * {children_count};
 
                                  let index = global_invocation_id.x;
                                  let total = arrayLength(&sort_indirection);
+
+                                 /*
+                                    let request_count_round_up =
+                                    ((request_count / {children_count}) + select(u32(0), u32(1), request_count % {children_count} != 0)) * {children_count};
                                  if(index * 2 >= request_count_round_up)
                                  {{ return; }}
+                                 */
 
                                  let sub_phase = (both_phase >> 16) & 0xFFFF;
                                  let phase = both_phase & 0xFFFF;
+                                 let phase_total_width = u32(1 << (phase + 1));
 
+                                 request_count = 0;
 
 
                                  // Check if phase is last
-                                 let phase_total_width = u32(1 << (phase + 1));
+                                 /*
                                  let phase_count = u32(ceil(log2(f32(request_count))));
                                  if(phase > phase_count)
                                  {{
-                                      //return;
+                                      return;
                                  }}
+                                 */
 
+                                 /*
                                  if(phase == 0)
                                  {{
                                     let a = index * 2;
@@ -242,6 +249,7 @@ where
                                     sort_indirection[a].child = a % {children_count};
                                     sort_indirection[b].child = b % {children_count};
                                  }}
+                                 */
 
                                  var swap_indices = vec2<u32>(0, 0);
 
@@ -255,7 +263,9 @@ where
                                  }}
 
                                  // Do swap
-                                 if(swap_indices.y >= request_count_round_up)
+
+                                 //if(swap_indices.y >= request_count_round_up)
+                                 if(swap_indices.y >= total)
                                  {{
                                     // Suppose that swap_indices.y is -inf, dont swap
                                     return;
@@ -306,7 +316,7 @@ where
                             @group(0) @binding(2) var<storage, read_write> request_counts: u32; 
 
                             @compute
-                            @workgroup_size(16)
+                            @workgroup_size(64)
                             fn main(
                                     @builtin(global_invocation_id) global_invocation_id: vec3<u32>
                                 )
@@ -398,7 +408,7 @@ where
                             @group(0) @binding(2) var<storage, read_write> request_counts: atomic<u32>; 
 
                             @compute
-                            @workgroup_size(16)
+                            @workgroup_size(64)
                             fn main(
                                     @builtin(global_invocation_id) global_invocation_id: vec3<u32>
                                 )
@@ -455,7 +465,7 @@ where
                             @group(0) @binding(2) var<storage, read_write> request_count: atomic<u32>; 
 
                             @compute
-                            @workgroup_size(16)
+                            @workgroup_size(64)
                             fn main(
                                     @builtin(global_invocation_id) global_invocation_id: vec3<u32>
                                 )
@@ -469,6 +479,10 @@ where
                                     {{
                                         for(var i = 0; i < {children_count}; i += 1)
                                         {{
+                                            if(request_buffer[index].children[i] != 0)
+                                            {{
+                                                atomicAdd(&request_count, 1);
+                                            }}
                                             request_buffer[index].children[i] = 0;
                                         }}
                                     }}
@@ -522,7 +536,7 @@ where
         compute_pass.set_pipeline(&self.reset_pipeline);
 
         let shader_invocations = self.cache_size; // one invocation per element
-        let workgroup_invocations = shader_invocations.div_ceil(16);
+        let workgroup_invocations = shader_invocations.div_ceil(64);
         compute_pass.dispatch_workgroups(workgroup_invocations as u32, 1, 1);
     }
 
@@ -535,8 +549,9 @@ where
 
         compute_pass.set_bind_group(0, Some(&self.bindgroup), &[]);
         let request_element_count = self.cache_size; // Each child slot is sorted
-        let workgroups_invocations = request_element_count.div_ceil(16);
+        let workgroups_invocations = request_element_count.div_ceil(64);
 
+        /*
         // = Perform list compaction
         
         // == Running sum
@@ -562,21 +577,22 @@ where
         // == Stream compaction
         compute_pass.set_pipeline(&self.compaction_pipeline);
         compute_pass.dispatch_workgroups(workgroups_invocations as u32, 1, 1);
+        */
 
         // == Bitonic sort
         compute_pass.set_pipeline(&self.sort_pipeline);
-        let sort_element_count = request_element_count * N * N * N;
+        let sort_element_count = self.cache_size * N * N * N;
         let phases_upper_bound = sort_element_count.next_power_of_two().ilog2();
 
         // Phase 0 dispatch all
-        for i in 0..phases_upper_bound
+        for i in 0..=phases_upper_bound
         {
             for j in 0..=i
             {
                 compute_pass.set_immediates(0, bytemuck::bytes_of(
                         &(i | (j << 16))
                         ));
-                compute_pass.dispatch_workgroups(sort_element_count.div_ceil(2).div_ceil(16) as u32, 1, 1);
+                compute_pass.dispatch_workgroups(sort_element_count.div_ceil(2).div_ceil(64) as u32, 1, 1);
             }
         }
     }
@@ -714,7 +730,7 @@ impl UsageBuffer
 
 
                             @compute
-                            @workgroup_size(16)
+                            @workgroup_size(64)
                             fn main(
                                     @builtin(global_invocation_id) global_invocation_id: vec3<u32>
                                 )
@@ -810,9 +826,9 @@ impl UsageBuffer
         let shader_invocations = self.cache_size.div_ceil(2); // bitonic sorting :
         // half as many shaders
         // per element
-        let workgroup_invocations = shader_invocations.div_ceil(16);
+        let workgroup_invocations = shader_invocations.div_ceil(64);
 
-        for i in 0..sort_steps
+        for i in 0..=sort_steps
         {
             for j in 0..=i
             {
@@ -842,8 +858,8 @@ impl StructureTable
         let pointer_table = device.create_buffer(&wgpu::BufferDescriptor {
             label: Some("structure_table_pointer_table"),
             size: size_of::<u32>() as u64,
-            usage: BufferUsages::STORAGE | BufferUsages::COPY_DST,
-            mapped_at_creation: false,
+            usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC,
+             mapped_at_creation: false,
         });
 
         StructureTable {
@@ -867,7 +883,7 @@ impl StructureTable
         let pointer_table = self.device.create_buffer(&wgpu::BufferDescriptor {
             label: Some("structure_table_pointer_table"),
             size: size_of::<u32>() as u64 * self.allocation_table.len() as u64,
-            usage: BufferUsages::STORAGE,
+            usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC | BufferUsages::COPY_DST,
             mapped_at_creation: false,
         });
 
@@ -986,11 +1002,25 @@ pub struct DestinationElement
     pub child: u32,
 }
 
+pub struct ColorBytes(pub u8, pub u8, pub u8, pub u8);
+
+impl From<Color> for ColorBytes
+{
+    fn from(value: Color) -> Self {
+        Self(
+            (value.0 * 255.) as u8,
+             (value.1 * 255.) as u8,
+              (value.2 * 255.) as u8,
+            (value.3 * 255.) as u8,
+        )
+    }
+}
+
 pub struct ColorPoolElement<const N: usize>
 where
     [(); N * N * N]:,
 {
-    colors: [Color; N * N * N],
+    colors: [ColorBytes; N * N * N],
 }
 
 pub struct LocationPoolElement
@@ -1343,7 +1373,7 @@ where
 
                 struct ColorPoolElement
                 {{
-                      colors: array<vec4<f32>, {children_count}>
+                      colors: array<u32, {children_count}>
                 }}
 
                 struct LocationPoolElement
@@ -1427,7 +1457,7 @@ where
                             @group(2) @binding(0) var<storage, read_write> requests: array<CacheNodeRequest>; 
 
                             @compute
-                            @workgroup_size(16)
+                            @workgroup_size(64)
                             fn main(
                                     @builtin(global_invocation_id) global_invocation_id: vec3<u32>
                                 )
@@ -1435,7 +1465,7 @@ where
                                  // One shader invocation per invocation on both structure table domain
                                  // and cache domain
                                  var index = global_invocation_id.x;
-                                 let max_requests_count = arrayLength(&requests);
+                                 let max_requests_count = min(arrayLength(&requests), arrayLength(&lru_list));
                                  total_request_count = min(max_requests_count, pools_request_count + structure_table_request_count);
                                  if(index >= total_request_count)
                                  {{
@@ -1586,7 +1616,7 @@ where
                             @group(2) @binding(3) var<storage, read> destinations: array<DestinationElement>; 
 
                             @compute
-                            @workgroup_size(16)
+                            @workgroup_size(64)
                             fn main(
                                     @builtin(global_invocation_id) global_invocation_id: vec3<u32>
                                 )
@@ -1594,7 +1624,8 @@ where
                                     // Copy with indirection
                                     let index = global_invocation_id.x;
                                     let total = arrayLength(&structure_nodes);
-                                    if(index >= total)
+                                    let total_cache = arrayLength(&lru_list);
+                                    if(index >= total || index >= total_cache)
                                     {{
                                         return;
                                     }}
@@ -1604,7 +1635,10 @@ where
                                     {{
                                         // Phase 1
                                         // Copy into cache page
-                                        structure_pool[overwritten_element] = structure_nodes[index];
+                                        for(var i = 0; i < {children_count}; i++)
+                                        {{
+                                            structure_pool[overwritten_element].pointers[i] = select(u32(0), u32(1)<<31, structure_nodes[index].pointers[i] != 0);
+                                        }}
                                         color_pool[overwritten_element] = color_nodes[index];
                                         location_pool[overwritten_element] = locations[index];
 
@@ -1620,7 +1654,7 @@ where
                                         if(destinations[index].child == 0xFFFFFFFF)
                                         {{
                                             structure_table_pointers[destinations[index].node] = new_pointer;
-                                        }}else
+                                        }}else if usage_buffer[destinations[index].node] != parameters.frame_timestamp + 1
                                         {{
                                             structure_pool[destinations[index].node].pointers[destinations[index].child] = new_pointer;
                                         }}
@@ -1662,7 +1696,7 @@ where
                             var<immediate> frame_timestamp: u32;
 
                             @compute
-                            @workgroup_size(16)
+                            @workgroup_size(64)
                             fn main(
                                     @builtin(global_invocation_id) global_invocation_id: vec3<u32>
                                 )
@@ -1863,7 +1897,7 @@ where
 
         // Compute necessary shader invocations
         let shader_invocation_count = request_target.size() / size_of::<CacheNodeRequest>() as u64;
-        let workgroup_invocation_count = shader_invocation_count.div_ceil(16);
+        let workgroup_invocation_count = shader_invocation_count.div_ceil(64);
         write_requests_pass.dispatch_workgroups(workgroup_invocation_count as u32, 1, 1);
     }
 
@@ -1930,7 +1964,7 @@ where
             // Compute dispatch amounts
             let shader_invocation_count =
                 insertion.structure_nodes.size() as usize / size_of::<StructurePoolElement<N>>();
-            let workgroup_invocations = shader_invocation_count.div_ceil(16);
+            let workgroup_invocations = shader_invocation_count.div_ceil(64);
 
             cache_insertion_pass.dispatch_workgroups(workgroup_invocations as u32, 1, 1);
         }
@@ -1944,7 +1978,7 @@ where
 
             // Compute dispatch amounts
             let shader_invocation_count = self.size + self.structure_table.allocation_table.len();
-            let workgroup_invocations = shader_invocation_count.div_ceil(16);
+            let workgroup_invocations = shader_invocation_count.div_ceil(64);
 
             cache_insertion_pass.dispatch_workgroups(workgroup_invocations as u32, 1, 1);
         }
@@ -1964,7 +1998,7 @@ where
             // Compute dispatch amounts
             let shader_invocation_count =
                 insertion.structure_nodes.size() as usize / size_of::<StructurePoolElement<N>>();
-            let workgroup_invocations = shader_invocation_count.div_ceil(16);
+            let workgroup_invocations = shader_invocation_count.div_ceil(64);
 
             cache_insertion_pass.dispatch_workgroups(workgroup_invocations as u32, 1, 1);
         }