starting compaction stuff

.gitattributes

@@ -0,0 +1,3 @@
+img.jpg filter=lfs diff=lfs merge=lfs -text
+img_low.jpg filter=lfs diff=lfs merge=lfs -text
+vxls_height.tif filter=lfs diff=lfs merge=lfs -text

.gitignore

@@ -1,6 +1,7 @@
 /target
 Cargo.lock
 
+imgs.tar.gz
 
 # Added by cargo
 #

Cargo.toml

@@ -14,10 +14,11 @@ env_logger = "0.11.11"
 fastapprox = "0.3.1"
 glam = "0.33.5"
 image = "0.25.10"
+indicatif = "0.18.6"
 itertools = "0.15.0"
 pollster = "1.0.1"
 rand = "0.10.2"
 rayon = "1.12.0"
 tiff = "0.11.3"
-wgpu = "30"
+wgpu = {version = "30", features = ["spirv"]}
 winit = "0.30.13"

shaders/Makefile

@@ -0,0 +1,4 @@
+all: voxel.spv
+
+%.spv: %.slang
+	slangc $< -O3 -fvk-use-entrypoint-name -target spirv -o $@

shaders/compaction.slang

@@ -0,0 +1,124 @@
+
+[[vk::binding(0, 0)]]
+RWStructuredBuffer<uint32_t> count_buffer;
+
+[[vk::binding(0, 1)]]
+RWStructuredBuffer<uint32_t> reduced_buffer;
+[[vk::binding(1, 1)]]
+RWStructuredBuffer<uint32_t> sum_buffer;
+[[vk::binding(2, 1)]]
+RWStructuredBuffer<uint32_t> compaction_buffer;
+
+groupshared uint32_t local_data[256 * 2];
+static uint32_t      THREAD_WIDTH = 256;
+static uint32_t      DATA_WIDTH = THREAD_WIDTH * 2;
+
+[numthreads(256, 1, 1)]
+[shader("compute")]
+void block_sum(
+    uint32_t3 workgroup_id: SV_GroupID,
+    uint32_t3 local_thread_id: SV_GroupThreadID,
+    uint32_t3 global_thread_id: SV_DispatchThreadID)
+{
+    // Perform sum in current block
+
+    // Copy local_datainto LDS with predicate
+    let thread_index = global_thread_id.x;
+    let local_thread_index = local_thread_id.x;
+    let total = count_buffer.getCount();
+
+    if (thread_index * 2 < total)
+    {
+        local_data[local_thread_index * 2] = select(count_buffer[thread_index * 2] != 0, 1, 0);
+    }
+    else
+    {
+        local_data[local_thread_index * 2] = 0;
+    }
+
+    if (thread_index * 2 + 1 < total)
+    {
+        local_data[local_thread_index * 2 + 1] = select(count_buffer[thread_index * 2 + 1] != 0, 1, 0);
+    }
+    else
+    {
+        local_data[local_thread_index * 2 + 1] = 0;
+    }
+
+    GroupMemoryBarrierWithGroupSync();
+
+    var width : uint32_t = 2;
+    while (width <= DATA_WIDTH)
+    {
+        let dest_index = width * (thread_index + 1) - 1;
+        let get_index = dest_index - (width / 2);
+        // println!("{}, {}", get_index, dest_index);
+        if (dest_index < DATA_WIDTH)
+        {
+            local_data[dest_index] += local_data[get_index];
+        }
+        width *= 2;
+        GroupMemoryBarrierWithGroupSync();
+    }
+
+    local_data[DATA_WIDTH - 1] = 0;
+    while (width >= 2)
+    {
+        let dest_index = width * (thread_index + 1) - 1;
+        let get_index = dest_index - (width / 2);
+        // println!("{}, {}", get_index, dest_index);
+        if (dest_index < DATA_WIDTH)
+        {
+            let self_data = local_data[dest_index];
+            local_data[dest_index] += local_data[get_index];
+            local_data[get_index] = self_data;
+        }
+        width /= 2;
+        GroupMemoryBarrierWithGroupSync();
+    }
+
+    // Block now contains running local sum
+    // Dump back to sum buffer
+    sum_buffer[2 * thread_index] = local_data[2 * local_thread_index];
+    sum_buffer[2 * thread_index + 1] = local_data[2 * local_thread_index + 1];
+
+    // Write to reduced buffer
+    reduced_buffer[workgroup_id.x] = local_data[DATA_WIDTH - 1];
+}
+
+[numthreads(1, 1, 1)]
+[shader("compute")]
+void linear_reduced_sum(
+    uint32_t3 workgroup_id: SV_GroupID,
+    uint32_t3 local_thread_id: SV_GroupThreadID,
+    uint32_t3 global_thread_id: SV_DispatchThreadID)
+{
+    let size = reduced_buffer.getCount();
+
+    // Perform exclusive sum
+    var running_sum : uint32_t = 0;
+    for (uint32_t i = 0; i < size; i++)
+    {
+        let value = reduced_buffer[i];
+        reduced_buffer[i] = running_sum;
+        running_sum += value;
+    }
+}
+
+[numthreads(256, 1, 1)]
+[shader("compute")]
+void uniform_add(
+    uint32_t3 workgroup_id: SV_GroupID,
+    uint32_t3 local_thread_id: SV_GroupThreadID,
+    uint32_t3 global_thread_id: SV_DispatchThreadID)
+{
+    let thread_index = global_thread_id.x;
+    let local_thread_index = local_thread_id.x;
+
+    // Gather
+    let reduced_value = reduced_buffer[global_thread_id.x];
+    // Apply
+    sum_buffer[thread_index * 2] += reduced_value;
+    sum_buffer[thread_index * 2 + 1] += reduced_value;
+}
+

shaders/voxel.slang

@@ -0,0 +1,471 @@
+struct PushConstants
+{
+    float4x4 view_proj;
+    float3   cam_pos;
+    uint32_t frame_timestamp;
+    uint32_t width;
+    uint32_t height;
+
+    uint32_t chunk_width;
+    uint32_t chunk_height;
+    uint32_t chunk_alt;
+}
+
+public struct VertexOutput
+{
+    public float4 position : SV_Position;
+
+    [vk::location(0)]
+    public float3 world_position;
+
+    [vk::location(1)]
+    public nointerpolation uint32_t structure_id;
+
+    [vk::location(2)]
+    public float3 cam_position;
+
+    [vk::location(3)]
+    public float3 chunk_position;
+}
+
+[[vk::push_constant]]
+uniform PushConstants constants;
+
+[shader("vertex")]
+VertexOutput chunk(
+    uint                     index: SV_VulkanVertexID,
+    [vk::location(0)] float3 chunk_position,
+    [vk::location(1)] uint   id)
+{
+    let cube_vertices : float3[8] =
+                            float3[](
+                                float3(0., 0., 0.),
+                                float3(0., 0., 1.),
+                                float3(1., 0., 1.),
+                                float3(1., 0., 0.),
+
+                                float3(0., 1., 0.),
+                                float3(0., 1., 1.),
+                                float3(1., 1., 1.),
+                                float3(1., 1., 0.), );
+    // clang-format off
+    let cube_faces: int[24] = int[](
+        // 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: int[6] = int[](
+        0, 1, 2, 1, 3, 2
+    );
+
+
+    let vertex = cube_vertices[cube_faces[quad_index * 4 + triangle_map[triangle_index]]];
+    let output_vertex = mul(constants.view_proj, float4(vertex + chunk_position, 1.0f));
+
+    VertexOutput vertex_output;
+    vertex_output.position = output_vertex;
+    vertex_output.world_position = vertex + chunk_position;
+    vertex_output.structure_id = id;
+    vertex_output.cam_position = constants.cam_pos;
+    vertex_output.chunk_position = chunk_position;
+
+    return vertex_output;
+}
+
+struct StructurePointer
+{
+    uint32_t value;
+    bool subdivided()
+    {
+        return (this.value & 0x80000000) != 0;
+    }
+
+    bool pointer_valid()
+    {
+        return (this.value & 0x40000000) != 0;
+    }
+
+    bool subdivided_valid()
+    {
+        return (this.value & 0xC0000000) == 0xC0000000;
+    }
+
+    uint32_t pointer()
+    {
+        return this.value & 0x3FFFFFFF;
+    }
+}
+
+struct ByteColor
+{
+    uint32_t byte_color;
+    
+    property uint32_t byte_r {
+        get {return byte_color & 0xFF;}
+    }
+
+    property uint32_t byte_g {
+        get {return (byte_color >> 8) & 0xFF;}
+    }
+
+    property uint32_t byte_b {
+        get {return (byte_color >> 16) & 0xFF;}
+    }
+
+    property uint32_t byte_a {
+        get {return byte_color >> 24;}
+    }
+
+    property float4 float_color {
+        get {return float4(
+                float(byte_r) / 255.,
+                float(byte_g) / 255.,
+                float(byte_b) / 255.,
+                float(byte_a) / 255.
+        ); }
+    }
+}
+
+struct StructurePoolElement
+{
+    uint32_t occupancy_low;
+    uint32_t occupancy_high;
+    StructurePointer pointers[64];
+}
+
+struct RequestBufferElement
+{
+    Atomic<uint32_t> requests[64];
+}
+
+struct ColorPoolElement
+{
+    ByteColor colors[64];
+}
+
+struct LocationPoolElement
+{
+    uint32_t structure_id;
+    uint32_t structure_locator;
+}
+
+[[vk::binding(0, 0)]] RWStructuredBuffer<StructurePoolElement> structure_pool;
+[[vk::binding(1, 0)]] RWStructuredBuffer<ColorPoolElement> color_pool;
+[[vk::binding(2, 0)]] RWStructuredBuffer<LocationPoolElement> location_pool;
+[[vk::binding(3, 0)]] RWStructuredBuffer<RequestBufferElement> request_buffer;
+[[vk::binding(4, 0)]] RWStructuredBuffer<uint32_t> usage_buffer;
+[[vk::binding(5, 0)]] RWStructuredBuffer<StructurePointer> structure_table_pointer;
+[[vk::binding(6, 0)]] RWStructuredBuffer<Atomic<uint32_t>> structure_table_request_buffer;
+
+uint32_t3 get_children_pos(float3 position, uint32_t scale_exp)
+{
+    return (asuint(position) >> scale_exp) & 3;
+}
+
+uint32_t get_children_index(float3 position, uint32_t scale_exp)
+{
+    // Get mantissa bits for this scale exp an retain bits for the specific children
+    uint32_t3 cell_position = (asuint(position) >> scale_exp) & 3;
+    return cell_position.x + cell_position.y * 4 + cell_position.z * 4 * 4;
+}
+
+uint64_t get_child_mask(uint32_t low, uint32_t high)
+{
+    return ((uint64_t)high << 32) | (uint64_t)low;
+}
+
+float3 floor_scale(float3 position, uint32_t scale_exp)
+{
+    uint32_t mask = ~0u << scale_exp;
+    return asfloat(asuint(position) & mask);
+}
+
+struct HitInformation
+{
+    bool hit;
+    float3 hit_pos;
+    float4 color;
+}
+
+// struct NodeStack
+// {
+//     uint32_t node_stack[5];
+// }
+
+groupshared uint32_t stack[256 * 5];
+
+HitInformation ray_march(float3 ray_direction, float3 ray_origin, uint32_t root_id, float dist_offset, uint32_t stack_index)
+{
+    float fov_deg = 100. / 1920.;
+    float fov_rad = (float.getPi() * fov_deg) / 180.;
+    float cone_factor = tan(fov_rad / 2.) * 2; // Horizontal size of pixel
+
+    let st_pointer = structure_table_pointer[root_id];
+    if(!st_pointer.subdivided())
+    {
+        var hit: HitInformation;
+        hit.hit = false;
+        return hit;
+    }
+
+    if(!st_pointer.pointer_valid())
+    {
+        // Record request
+        structure_table_request_buffer[root_id].add(1);
+        var hit: HitInformation;
+        hit.hit = false;
+        return hit;
+    }
+
+    ray_origin += float3(1.);
+    ray_origin = clamp(ray_origin , float(1.), asfloat(0x3fffffff));
+    ray_origin = select(ray_direction > 0., asfloat(asuint(ray_origin) ^ 0x007fffff), ray_origin);
+    uint32_t child_mirror = 0;
+    if(ray_direction.x > 0.) {child_mirror |= 3;}
+    if(ray_direction.y > 0.) {child_mirror |= 3 << 2;}
+    if(ray_direction.z > 0.) {child_mirror |= 3 << 4;}
+    ray_direction = -abs(ray_direction);
+
+    float3 pos = ray_origin;
+
+    uint32_t scale_exp = 23 - 2;
+    uint32_t node_stack[5] = 
+    {
+        0
+    };
+
+    uint32_t current_node_index = structure_table_pointer[root_id].pointer();
+    usage_buffer[current_node_index] = constants.frame_timestamp;
+    stack[stack_index * 5 + 10 - scale_exp / 2] = current_node_index;
+    //node_stack[10 - scale_exp / 2] = current_node_index;
+
+    [loop]
+    for(uint32_t iter = 0; iter < 500; iter ++)
+    {
+        uint32_t child_index = get_children_index(pos, scale_exp) ^ child_mirror;
+        StructurePointer current_node = structure_pool[current_node_index].pointers[child_index];
+
+        // Scale computations
+        let cone_size = (length(ray_origin - pos) + dist_offset) * cone_factor;
+        let exponent = 
+            select(
+                cone_size == 0.,
+                0,
+                23 - (127 - (asuint(cone_size) >> 23))
+            );
+
+        while(
+            current_node.subdivided_valid() && 
+            scale_exp - 2 > exponent
+            )
+        {
+            scale_exp -= 2;
+            current_node_index = current_node.pointer();
+            stack[stack_index * 5 + 10 - scale_exp / 2] = current_node_index;
+            //node_stack[10 - scale_exp / 2] = current_node_index;
+            child_index = get_children_index(pos, scale_exp) ^ child_mirror;
+            current_node = structure_pool[current_node_index].pointers[child_index];
+            // Write usage
+            usage_buffer[current_node_index] = constants.frame_timestamp;
+        }
+
+        // Request subdiv
+        if(current_node.subdivided() && !current_node.pointer_valid())
+        {
+            request_buffer[current_node_index].requests[child_index].add(1);
+        }
+
+        if(color_pool[current_node_index].colors[child_index].byte_a != 0)
+        {
+            var hit: HitInformation;
+            hit.hit = true;
+            hit.hit_pos = pos - float3(1.);
+            hit.color = color_pool[current_node_index].colors[child_index].float_color;
+            return hit;
+        }
+
+        uint64_t occupancy = get_child_mask(structure_pool[current_node_index].occupancy_low, structure_pool[current_node_index].occupancy_high);
+        uint32_t adv_scale_exp = scale_exp;
+        if(((occupancy >> (child_index & 0b101010)) & 0x00330033) == 0)
+        {
+            adv_scale_exp ++;
+        }
+
+        // Perform dda
+        // Compute correct exponent, and shift it into the exponent part of floatt
+        let child_pos : float3 = floor_scale(pos, adv_scale_exp);
+        // Intersection t
+        let inter_ts : float3 = (child_pos - ray_origin) / ray_direction;
+        float inter_t = min(inter_ts.x, min(inter_ts.y, inter_ts.z));
+        //return float4(inter_t);
+
+        // Perform dda step
+        //let neighbor_max = asint(child_pos) + select(inter_t == inter_ts, -1, (1 << adv_scale_exp) - 1);
+        let neighbor_max = asint(child_pos) + select(inter_t == inter_ts, -1, (1 << adv_scale_exp) - 1);
+        pos = min(ray_origin + ray_direction * inter_t, asfloat(neighbor_max));
+
+        // Find most common ancestor
+        uint32_t3 diffs = asuint(child_pos) ^ asuint(pos);
+        uint32_t diff = (diffs.x | diffs.y | diffs.z);
+        
+        int32_t common_depth = (1 + (22 - firstbithigh(diff)) / 2) * 2;
+        if(common_depth <= 0)
+        {
+            break;
+        }
+
+        scale_exp = 23 - common_depth;
+        current_node_index = stack[stack_index * 5 + 10 - scale_exp / 2];
+        //current_node_index = node_stack[10 - scale_exp / 2];
+    }
+
+
+    var hit: HitInformation;
+    hit.hit = false;
+    return hit;
+}
+
+struct FragmentOutput
+{
+    float depth : SV_Depth;
+    float4 color : SV_Target<0>;
+}
+
+/*
+//[earlydepthstencil]
+[shader("fragment")]
+FragmentOutput fragment(VertexOutput vertex_out)
+{
+    let ray_direction = normalize(vertex_out.world_position - vertex_out.cam_position);
+    let intersection_t = box_intersect(vertex_out.cam_position, ray_direction, vertex_out.chunk_position, vertex_out.chunk_position + float3(1.));
+    let local_ray_origin = max(intersection_t.x, 0.) * ray_direction + vertex_out.cam_position - vertex_out.chunk_position;
+    // Figure out intersection
+    let hit = ray_march(ray_direction, local_ray_origin, vertex_out.structure_id, max(0., intersection_t.x));
+    let world_hit_pos = hit.hit_pos + vertex_out.chunk_position;
+    let clip = mul(constants.view_proj, float4(world_hit_pos, 1.));
+    let depth = clip.z / clip.w;
+
+    var frag_out : FragmentOutput;
+    frag_out.depth = depth;
+    frag_out.color = hit.color;
+
+    return frag_out;
+}
+*/
+
+bool3 min_mask(float3 val)
+{
+    let min_val = min(val.x, min(val.y, val.z));
+    return val == min_val;
+}
+
+[[vk::binding(0, 1)]] 
+[[format("rgba32f")]]
+WTexture2D<float4> output_texture;
+
+float3 get_ray_direction(uint32_t2 pixel_loc)
+{
+    let ndc_loc_x = (float)pixel_loc.x / (float)constants.width * 2. - 1.;
+    let ndc_loc_y = 1. - (float)pixel_loc.y / (float)constants.height * 2.;
+    var world_loc = mul(constants.view_proj, float4(ndc_loc_x, ndc_loc_y, 1., 1.));
+    world_loc /= world_loc.w;
+    return normalize(world_loc.xyz - constants.cam_pos);
+}
+
+[shader("compute")]
+[numthreads(16, 16, 1)]
+void ray_march_compute(uint32_t3 location : SV_DispatchThreadID, uint32_t3 local_location: SV_GroupThreadID)
+{
+    if(location.x >= constants.width || location.y >= constants.height)
+    {
+        return;
+    }
+
+    let stack_index = local_location.x + local_location.y * 16;
+    let ray_direction = get_ray_direction(location.xy);
+    var inter = box_intersect(constants.cam_pos, ray_direction, float3(0.), float3(constants.chunk_width, constants.chunk_alt, constants.chunk_height));
+
+    // Clear if out of box
+    if(inter.y <= inter.x || inter.y <= 0.)
+    {
+        output_texture.Store(location.xy, float4(0.));
+        return;
+    }
+    inter.x = max(0., inter.x);
+    // float3 position = constants.cam_pos + ray_direction * inter.x;
+    // int32_t3 current_voxel = clamp(
+    //     int32_t3(floor(position)), 
+    //     int32_t3(0),
+    //     int32_t3(constants.chunk_width - 1, constants.chunk_alt - 1, constants.chunk_height - 1)
+    // );
+    // let chunk_index = current_voxel.y + current_voxel.z * constants.chunk_alt + current_voxel.x * constants.chunk_alt * constants.chunk_height;
+    // let hit = ray_march(ray_direction, position - float3(current_voxel), chunk_index, length(position - constants.cam_pos), stack_index);
+    // if(hit.hit)
+    // {
+    //     output_texture.Store(location.xy, float4(hit.color));
+    // }
+
+
+    let start_position = inter.x * ray_direction + constants.cam_pos;
+    // FVT
+    int32_t3 current_voxel = clamp(
+        int32_t3(floor(start_position)), 
+        int32_t3(0),
+        int32_t3(constants.chunk_width - 1, constants.chunk_alt - 1, constants.chunk_height - 1)
+    );
+    //int32_t3 offset = int32_t3(sign(ray_direction));
+    //float3 delta = abs(1. / ray_direction);
+    float3 t = select(ray_direction > 0., current_voxel + int32_t3(1) - start_position, start_position - current_voxel) / abs(ray_direction);
+    t += inter.x;
+
+    [loop]
+    while(true)
+    {
+        {
+            let off = t - abs(1. / ray_direction);
+            let t_adv = max(max(off.x, max(off.y, off.z)), 0.);
+            let chunk_index = current_voxel.y + current_voxel.z * constants.chunk_alt + current_voxel.x * constants.chunk_alt * constants.chunk_height;
+            let position = constants.cam_pos + ray_direction * t_adv;
+            let hit = ray_march(ray_direction, position - float3(current_voxel), chunk_index, t_adv, stack_index);
+            if(hit.hit)
+            {
+                output_texture.Store(location.xy, hit.color);
+                return;
+            }
+        }
+
+        let min_mask = min_mask(t);    
+        t += select(min_mask, abs(1. / ray_direction), float3(0.));
+        current_voxel += select(min_mask, int32_t3(sign(ray_direction)), int32_t3(0));
+
+        if(any(current_voxel < 0) || any(current_voxel >= int32_t3(constants.chunk_width, constants.chunk_alt, constants.chunk_height)))
+        {
+            output_texture.Store(location.xy, float4(0.));
+            return;
+        }
+    }
+}
+
+float2 box_intersect(float3 origin, float3 ray_direction, float3 box_min, float3 box_max)
+{
+    let min_ts = (box_min - origin) / ray_direction;
+    let max_ts = (box_max - origin) / ray_direction;
+
+    let far_ts = max(min_ts, max_ts);
+    let near_ts = min(min_ts, max_ts);
+
+    let far_t = min(far_ts.x, min(far_ts.y, far_ts.z));
+    let near_t = max(near_ts.x, max(near_ts.y, near_ts.z));
+    return float2(near_t, far_t);
+}

shaders/voxel.temp.wgsl

@@ -0,0 +1,399 @@
+struct VertexOutput
+{
+    @builtin(position) postion: vec4<f32>,
+    @location(0) @interpolate(flat) chunk_index: u32,
+    @location(1) color: vec4<f32>,
+    @location(2) cam_pos: vec3<f32>,
+    @location(3) world_pos: vec3<f32>,
+    @location(4) @interpolate(flat) structure_id: u32,
+    @location(5) chunk_position: vec3<f32>
+}
+
+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<u32, 64>
+}
+
+struct LocationPoolElement
+{
+    structure_id: u32,
+    structure_locator: u32
+}
+
+struct SortedRequestsElement
+{
+    node: u32,
+    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>;
+@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>>;
+
+struct FragmentOutput {
+    @location(0) color: vec4<f32>,
+    @builtin(frag_depth) depth: f32, // Equivalent to gl_FragDepth
+}
+
+@vertex
+fn chunk(@builtin(vertex_index) index: u32, @location(0) position: vec3<f32>, @location(1) id: u32) -> @builtin(position) vec4<f32>
+{
+    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 + position, 1.0f);
+
+    return output_vertex;
+}
+
+
+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) & 1) != 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)));
+}
+
+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. / 1920.;
+    let fovy_rad = (fovy_deg * 3.14) / 180.;
+    let cone_factor = tan(fovy_rad / 2.) * 2.;
+
+    let st_pointer = structure_table_pointer[root_id];
+
+    
+    if (!node_subdivided(st_pointer))
+    {
+        discard;
+        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); 
+
+        discard;
+        var result: HitResult;
+        result.color = vec4(0., 1., 0., 1.);
+        result.hit_pos = ray_origin;
+        return result;
+    }
+    //var current_node = node_pointer(st_pointer);
+
+    var dfs_stack = array<u32, 6>(node_pointer(st_pointer), 0, 0, 0, 0, 0);
+    var current_depth = 0;
+    var current_node = dfs_stack[current_depth];
+
+    usage_buffer[current_node] = constants.frame_timestamp;
+
+    // Start location
+    //let voxel_dir = select(vec3(-1), vec3(1), ray_dir >= vec3(0.));
+    var node_shift = (max_depth - current_depth) * 2;
+
+    var child_size = 1 << u32(node_shift - 2);
+    var node_size = 1 << u32(node_shift);
+
+    var pos_origin = clamp(ray_origin * f32(1 << u32(max_depth * 2)), vec3(0.), vec3(f32(node_size) - 1.));
+    var voxel = vec3<i32>(pos_origin);
+    var far_t = 0.;
+    var inv_ray_dir = 1. / ray_dir;
+    var ray_positive = ray_dir > vec3(0.);
+    var step_dir = select(vec3(-1), vec3(1), ray_positive);
+
+    for(var iter = 0; iter < 400; iter ++)
+    {
+        // Shift into voxel position
+        node_shift = (max_depth - current_depth) * 2;
+        child_size = 1 << u32(node_shift - 2);
+
+        // Compute child position position from voxel position
+        var child_pos = (voxel >> vec3(u32(node_shift - 2))) & vec3(3);
+        // Compute child index in pointers
+        var child_index = child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4;
+        // Candidate child pointer
+        var pointer = structure_pool[current_node].pointers[child_index];
+        
+        // Descent loop
+        let min_child_size = (length(vec3<f32>(voxel) - pos_origin) + dist_offset_voxel) * cone_factor;
+        while(node_subdivided(pointer) && node_pointer_valid(pointer) &&
+            f32(child_size / 4) >= min_child_size
+            )
+        {
+
+            // Descend 
+            current_depth += 1;
+
+            // Try to descend again
+            node_shift = (max_depth - current_depth) * 2;
+            child_size = 1 << u32(node_shift - 2);
+            child_pos = (voxel >> vec3(u32(node_shift - 2))) & vec3(3);
+            current_node = node_pointer(pointer);
+            dfs_stack[current_depth] = current_node;
+            child_index = child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4;
+
+            pointer = structure_pool[current_node].pointers[child_index];
+
+            // Record usage in usage buffer
+            usage_buffer[current_node] = constants.frame_timestamp;
+        }
+
+        // If we could not descencd, request the child
+        if(node_subdivided(pointer) && !node_pointer_valid(pointer) && 
+            f32(child_size / 4) >= min_child_size)
+        {
+            // Record request
+            atomicAdd(&request_buffer[dfs_stack[current_depth]].requests[child_index], 1);
+        }
+
+        // Check color
+        let color = color_pool[current_node].colors[child_index];
+        if(((color >> 24) & 0xFF) != 0)
+        {
+            var result: HitResult;
+            result.color = unpack_color(color);
+            result.hit_pos = (far_t / f32(1 << u32(max_depth * 2))) * ray_dir + ray_origin;
+            return result;
+        }
+
+        // Advance
+        child_pos = voxel & vec3(i32(0xFFFFFFFF << u32(node_shift - 2)));
+        let far_wall = child_pos + select(vec3(0), vec3(child_size), ray_positive);
+        let far_wall_inter = (vec3<f32>(far_wall) - pos_origin) * inv_ray_dir;
+        far_t = min(min(far_wall_inter.x, far_wall_inter.y), far_wall_inter.z);
+
+        // Perform dda step on the children scale
+        let next_child = select(child_pos, child_pos + step_dir * vec3(child_size), vec3(far_t) == far_wall_inter);
+        
+        let previous_voxel = voxel;
+        voxel = clamp(vec3<i32>(pos_origin + far_t * ray_dir), next_child, next_child + vec3(child_size) - vec3(1));
+
+        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;
+}
+
+@fragment
+fn fragment() -> @location(0) vec4<f32>
+{
+    return vec4(1., 0., 0., 1.) ;
+}
+
+@early_depth_test(less_equal)
+@fragment
+fn _fragment(in: VertexOutput) -> FragmentOutput
+{
+    //frag_out.color = vec4<f32>(2 * 0.01 / (100. + 0.01 - depth * (100. - 0.01)));
+    let ray_dir = normalize(in.world_pos - in.cam_pos);
+    let interp = box_inter(in.cam_pos - in.chunk_position, ray_dir, vec3(0.), vec3(1));
+    let ray_origin = (in.cam_pos - in.chunk_position) + ray_dir * (max(0., interp.x));
+
+
+    let result = new_traverse(ray_dir, ray_origin, in.structure_id, length(in.cam_pos - (ray_origin + in.chunk_position)));
+    let clip_pos = constants.view_proj * vec4(result.hit_pos + in.chunk_position, 1.);
+    let depth = clip_pos.z / clip_pos.w;
+    var frag_out: FragmentOutput;
+    //frag_out.color = result.color;
+    frag_out.color = result.color;
+    frag_out.depth = depth;
+    return frag_out;
+
+    //return vec4<f32>(ray_origin, 1.);
+    //return frag_out;
+    //return vec4(interp.y / 10.);
+}
+
+/*
+@fragment
+fn fragment(in: VertexOutput) -> @location(0) vec4<f32>
+{
+    let st = structure_table_pointer[0];
+    let subdivided = ((st >> 31) & 1) != 0;
+    let pointer_valid = ((st >> 30) & 1) != 0;
+    // Request stuff
+    atomicAdd(&structure_table_request_buffer[0], 1); 
+    if(subdivided && !pointer_valid)
+    {
+        return vec4(0., 1., 0., 1.);
+    }
+    return vec4(1., 0., 0., 1.);
+}
+*/
+

shaders/voxel.wgsl

@@ -216,14 +216,16 @@ fn new_traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_id: u32, dist_of
 {
     let max_depth = 5;
     let dist_offset_voxel = dist_offset * f32(1 << u32(max_depth * 2));
-    let fovy_deg = 100.;
-    let cone_factor = tan((fovy_deg / 180.) * 3.14159) * 2.;
+    let fovy_deg = 100. / 1920.;
+    let fovy_rad = (fovy_deg * 3.14) / 180.;
+    let cone_factor = tan(fovy_rad / 2.) * 2.;
 
     let st_pointer = structure_table_pointer[root_id];
 
     
     if (!node_subdivided(st_pointer))
     {
+        discard;
         var result: HitResult;
         result.color = vec4(0., 1., 0., 1.);
         result.hit_pos = ray_origin;
@@ -235,6 +237,7 @@ fn new_traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_id: u32, dist_of
         // Send request on structure table
         atomicAdd(&structure_table_request_buffer[root_id], 1); 
 
+        discard;
         var result: HitResult;
         result.color = vec4(0., 1., 0., 1.);
         result.hit_pos = ray_origin;
@@ -242,57 +245,67 @@ fn new_traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_id: u32, dist_of
     }
     //var current_node = node_pointer(st_pointer);
 
-    // Record usage
     var dfs_stack = array<u32, 6>(node_pointer(st_pointer), 0, 0, 0, 0, 0);
     var current_depth = 0;
+    var current_node = dfs_stack[current_depth];
 
-    usage_buffer[dfs_stack[current_depth]] = constants.frame_timestamp;
+    usage_buffer[current_node] = constants.frame_timestamp;
 
     // Start location
     //let voxel_dir = select(vec3(-1), vec3(1), ray_dir >= vec3(0.));
-    var node_size = 1 << u32(((max_depth - current_depth) * 2));
-    var child_size = node_size / 4;
+    var node_shift = (max_depth - current_depth) * 2;
+
+    var child_size = 1 << u32(node_shift - 2);
+    var node_size = 1 << u32(node_shift);
+
     var pos_origin = clamp(ray_origin * f32(1 << u32(max_depth * 2)), vec3(0.), vec3(f32(node_size) - 1.));
     var voxel = vec3<i32>(pos_origin);
     var far_t = 0.;
+    var inv_ray_dir = 1. / ray_dir;
+    var ray_positive = ray_dir > vec3(0.);
+    var step_dir = select(vec3(-1), vec3(1), ray_positive);
 
     for(var iter = 0; iter < 400; iter ++)
     {
         // Compute child position
-        node_size = 1 << u32(((max_depth - current_depth) * 2));
-        child_size = node_size / 4;
-        var child_pos = (voxel / child_size) % 4;
-        var pointer = structure_pool[dfs_stack[current_depth]].pointers[child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4];
+        node_shift = (max_depth - current_depth) * 2;
+        child_size = 1 << u32(node_shift - 2);
 
+        var child_pos = (voxel >> vec3(u32(node_shift - 2))) & vec3(3);
+        var child_index = child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4;
+        var pointer = structure_pool[current_node].pointers[child_index];
+        
+        let min_child_size = (length(vec3<f32>(voxel) - pos_origin) + dist_offset_voxel) * cone_factor;
         while(node_subdivided(pointer) && 
-            !((length(vec3<f32>(voxel) - pos_origin) + dist_offset) * cone_factor >= f32(node_size / 4))
+            f32(child_size / 4) >= min_child_size
             )
         {
-
             if(!node_pointer_valid(pointer) && node_subdivided(pointer))
             {
                 // Record request
-                atomicAdd(&request_buffer[dfs_stack[current_depth]].requests[child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4], 1);
+                atomicAdd(&request_buffer[dfs_stack[current_depth]].requests[child_index], 1);
                 break;
             }
 
             // Descend 
             current_depth += 1;
 
-            node_size /= 4;
-            child_size /= 4;
-            child_pos = (voxel / child_size) % 4;
-            dfs_stack[current_depth] = node_pointer(pointer);
+            node_shift = (max_depth - current_depth) * 2;
+            child_size = 1 << u32(node_shift - 2);
+            child_pos = (voxel >> vec3(u32(node_shift - 2))) & vec3(3);
+            current_node = node_pointer(pointer);
+            dfs_stack[current_depth] = current_node;
+            child_index = child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4;
 
-            pointer = structure_pool[dfs_stack[current_depth]].pointers[child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4];
+            pointer = structure_pool[current_node].pointers[child_index];
 
             // Record usage
-            usage_buffer[dfs_stack[current_depth]] = constants.frame_timestamp;
+            usage_buffer[current_node] = constants.frame_timestamp;
         }
 
 
         // Check color
-        let color = color_pool[dfs_stack[current_depth]].colors[child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4];
+        let color = color_pool[current_node].colors[child_index];
         if(((color >> 24) & 0xFF) != 0)
         {
             var result: HitResult;
@@ -302,13 +315,14 @@ fn new_traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_id: u32, dist_of
         }
 
         // Advance
-        child_pos = (voxel / child_size) * child_size;
-        let far_wall = child_pos + select(vec3(0), vec3(child_size), ray_dir > vec3(0.));
-        let far_wall_inter = (vec3<f32>(far_wall) - pos_origin) / ray_dir;
+        child_pos = voxel & vec3(i32(0xFFFFFFFF << u32(node_shift - 2)));
+        let far_wall = child_pos + select(vec3(0), vec3(child_size), ray_positive);
+        let far_wall_inter = (vec3<f32>(far_wall) - pos_origin) * inv_ray_dir;
         far_t = min(min(far_wall_inter.x, far_wall_inter.y), far_wall_inter.z);
 
         // Perform dda step on the children scale
-        let next_child = select(child_pos, child_pos + select(vec3(-1), vec3(1), ray_dir > vec3(0.)) * vec3(child_size), vec3(far_t) == far_wall_inter);
+        //let next_child = select(child_pos, child_pos + select(vec3(-1), vec3(1), ray_dir > vec3(0.)) * vec3(child_size), vec3(far_t) == far_wall_inter);
+        let next_child = select(child_pos, child_pos + step_dir * vec3(child_size), vec3(far_t) == far_wall_inter);
         
         let previous_voxel = voxel;
         voxel = clamp(vec3<i32>(pos_origin + far_t * ray_dir), next_child, next_child + vec3(child_size) - vec3(1));
@@ -331,7 +345,7 @@ fn new_traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_id: u32, dist_of
         let common_depth = ((countLeadingZeros(bit_diffs_lowest) - i32(32 - max_depth * 2)) / 2);
 
         current_depth = common_depth;
-        //current_node = dfs_stack[current_depth];
+        current_node = dfs_stack[current_depth];
     }
 
     // Iter max color
@@ -341,204 +355,6 @@ fn new_traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_id: u32, dist_of
     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 * f32(node_size), vec3(0.), vec3(f32(node_size) - 1.));
-    var voxel = vec3<i32>(pos_origin);
-    var last_voxel = voxel;
-
-    let wall_offset = select(vec3(0), vec3(1), ray_dir > vec3(0.));
-
-    let max_depth = u32(5);
-    var adaptive_depth = i32(max_depth);
-    var far_t = 0.;
-
-    let ray_dir_inv = 1. / ray_dir;
-    let fma_offset = - pos_origin * ray_dir_inv;
-
-    //let depth_limit = 3;
-    for(var iter = 0; iter < 400; iter ++)
-    {
-
-        // Our ray is currently touching a voxel.
-        // Descend to the lowest node that contains this voxel
-
-        // Position of the child we are in
-        var child_pos = (vec3<u32>(voxel) >> vec3<u32>((max_depth - u32(current_depth + 1)) * 2)) & vec3<u32>(3); // Hardcode for 4-tree
-        var child_index = child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4;
-
-        // Current node has been used. report
-        usage_buffer[current_node] = constants.frame_timestamp;
-        
-        while(
-            node_subdivided(structure_pool[current_node].pointers[child_index]) &&
-            (local_dist_offset + far_t) * cone_size_factor < f32(node_size_lut[current_depth])
-            )
-        {
-            if(!node_pointer_valid(structure_pool[current_node].pointers[child_index]))
-            {
-                atomicAdd(&request_buffer[current_node].requests[child_index], 1);
-                break;
-            }
-            // Child node is subdivided, we go in, save position in stack
-            current_node = node_pointer(structure_pool[current_node].pointers[child_index]);
-            usage_buffer[current_node] = constants.frame_timestamp;
-            current_depth += 1;
-            dfs_stack[current_depth] = current_node;
-            node_size = node_size_lut[current_depth];
-
-            child_pos = (vec3<u32>(voxel) >> vec3<u32>((max_depth - u32(current_depth + 1)) * 2)) & vec3<u32>(3); // Hardcode for 4-tree
-            child_index = child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4;
-            child_size = node_size / 4;
-        }
-        usage_buffer[current_node] = constants.frame_timestamp;
-
-
-
-
-        // At this point current_depth is the depth of the node that contains the voxel
-        // child_pos and child_index relate to the specific child of the node that contains this voxel
-        // It is guaranteed that the child is leave
-
-        // Check current leave's color
-        let color = unpack_color(color_pool[current_node].colors[child_index]);
-        if(color.w != 0.) // Not transparent
-        {
-            /*
-            let k = child_pos.x + child_pos.y + child_pos.z;
-            let w = voxel.x + voxel.y + voxel.z;
-            let x = select(0.5, 1., k % 2 == 0) * select(0.8, 1., w % 2 == 0);
-
-            var div = 1;
-            var overlay = 1.;
-            for(var i = 1; i <= 5; i++)
-            {
-                let x = (voxel.x / div + voxel.y / div + voxel.z / div) % 2 == 0;
-                overlay -= select(0., 1. / (f32(i) * 2.5), x);
-                div *= 4;
-            }
-            */
-
-            return color;
-            //return overlay * color;
-        }
-
-        // Voxel and whole child containing it is empty
-        
-        // Perform a step through the children of the node
-        let child_position = (voxel / child_size) * child_size;
-        let far_corner = child_position + wall_offset * child_size;
-        //let far_ts = (vec3<f32>(far_corner) - pos_origin) / ray_dir; // TODO: Turn into fma
-        let far_ts = fma(vec3<f32>(far_corner), ray_dir_inv, fma_offset);
-        far_t = min(min(far_ts.x, far_ts.y), far_ts.z);
-
-        let next_child_min = select(child_position, child_position + voxel_dir * child_size, vec3(far_t) == far_ts);
-        let next_child_max = next_child_min + vec3(child_size) - vec3(1);
-        
-        // The ray (far_t) is now touching the new child to explore
-        // Find out which actual voxel we are touching
-        let previous_voxel = voxel;
-        let float_voxel = clamp(vec3<i32>(pos_origin + far_t * ray_dir), next_child_min, next_child_max);
-        /*
-        voxel = vec3<i32>(
-            floor(
-            select(
-                float_voxel - vec3(0.5),
-                float_voxel + vec3(0.5),
-                ray_dir > vec3(0.)
-            ))
-        );
-        */
-        //voxel = vec3<i32>(round(float_voxel));
-        //voxel = voxel_from_wall(float_voxel, ray_dir);
-        //voxel = voxel_from_wall(float_voxel, ray_dir);
-        voxel = float_voxel;
-        if(any(voxel < vec3(0)) || any(voxel >= vec3(node_size_lut[0])))
-        {
-            //return vec4(f32(iter) / 100.);
-            discard;
-        }
-
-        // We touched a voxel as if we explored blocks sized by the child size of the current node.
-        // But we might have exited the current node.
-        
-        // If this is the case we have to walk back up the tree
-        // And then back down to the next node over
-
-        // As such we find the lowest ancestor that can contain both the privous voxel (in node) and the new voxel (out of node)
-        let bit_diffs = voxel ^ previous_voxel;
-        let bit_diffs_lowest = bit_diffs.x | bit_diffs.y | bit_diffs.z;
-
-        let flb = ((countLeadingZeros(bit_diffs_lowest) - i32(32 - max_depth * 2)) / 2);
-        let common_depth = flb;
-
-        current_depth = common_depth;
-        node_size = node_size_lut[current_depth];
-        child_size = node_size / 4;
-        current_node = dfs_stack[current_depth];
-        
-        // Figure out current voxel position
-        //voxel = vec3<i32>(ray_origin + ray_dir * t); 
-
-    }
-    return vec4<f32>(1., 0., 1., 1.);
-
-}
-
 @early_depth_test(less_equal)
 @fragment
 fn fragment(in: VertexOutput) -> FragmentOutput
@@ -549,10 +365,11 @@ fn fragment(in: VertexOutput) -> FragmentOutput
     let ray_origin = (in.cam_pos - in.chunk_position) + ray_dir * (max(0., interp.x));
 
 
-    let result = new_traverse(ray_dir, ray_origin, in.structure_id, length(in.cam_pos - ray_origin));
+    let result = new_traverse(ray_dir, ray_origin, in.structure_id, length(in.cam_pos - (ray_origin + in.chunk_position)));
     let clip_pos = constants.view_proj * vec4(result.hit_pos + in.chunk_position, 1.);
     let depth = clip_pos.z / clip_pos.w;
     var frag_out: FragmentOutput;
+    //frag_out.color = result.color;
     frag_out.color = result.color;
     frag_out.depth = depth;
     return frag_out;

src/host_production.rs

@@ -0,0 +1,15 @@
+// Contains useful facilities to render data streamed in from the host
+
+// Can produce a voxel given
+// - Its depth
+// - Its position within the chunk
+// - The chunks position
+pub trait ChunkVoxelProducer
+{
+    fn produce_voxel(
+        &mut self,
+        depth: usize,
+        chunk_position: (usize, usize, usize),
+        voxel_position: (usize, usize, usize),
+    );
+}

src/main.rs

@@ -1,63 +1,47 @@
 #![feature(generic_const_exprs)]
+#![feature(float_algebraic)]
 
-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::ops::Div;
 use std::sync::Arc;
-use std::sync::atomic::AtomicBool;
 use std::sync::mpsc::sync_channel;
 
 use bytemuck::Pod;
 use bytemuck::Zeroable;
+use bytemuck::cast_slice;
 use crevice::std140::AsStd140;
-use crevice::std430::AsStd430;
 use egui::Color32;
-use egui::Label;
-use egui::emath::fast_midpoint;
 use egui::mutex::Mutex;
 use egui_plot::BarChart;
 use glam::Mat4;
 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;
-use wgpu::BindGroupLayoutEntry;
+use wgpu::BindGroupLayout;
 use wgpu::Buffer;
 use wgpu::BufferUsages;
-use wgpu::DepthBiasState;
+use wgpu::ComputePipeline;
 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::Origin3d;
 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::include_spirv;
+use wgpu::include_wgsl;
 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::MouseScrollDelta;
@@ -73,32 +57,25 @@ 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;
-use crate::voxel::cache::UsageBuffer;
-use crate::voxel::cache::VoxelCache;
-use crate::voxel::gpu::ExplicitNTreeNode;
-use crate::voxel::pipeline::ChunkHandle;
-use crate::voxel::pipeline::ChunkObject;
-use crate::voxel::pipeline::VoxelPipeline;
-use crate::voxel::sparse::Color;
-use crate::voxel::sparse::NTree;
-use crate::voxel::sparse::NTreeNodeLocator;
+use crate::sparse_tree::NTreeNodeLocator;
+use crate::voxel_cache::VoxelCache;
+use crate::voxel_cache::data::CacheNodeRequest;
+use crate::voxel_cache::data::CacheResponse;
+use crate::voxel_cache::data::ColorBytes;
+use crate::voxel_cache::data::DestinationElement;
+use crate::voxel_cache::data::LocationPoolElement;
+use crate::voxel_cache::data::StructurePoolElement;
+use crate::voxel_cache::producer_interface::CacheProducerInterface;
+use crate::voxel_cache::producer_interface::CacheRequest;
 
 mod camera;
 mod egui_renderer;
 mod producers;
-mod voxel;
-//mod tree;
+mod sparse_tree;
+mod voxel_cache;
 //
 
 struct State
@@ -110,16 +87,22 @@ struct State
     size: winit::dpi::PhysicalSize<u32>,
     surface: wgpu::Surface<'static>,
     depth_buffer: (wgpu::Texture, wgpu::TextureView),
+    target_texture: (wgpu::Texture, wgpu::TextureView),
     surface_format: wgpu::TextureFormat,
+    target_blitter: wgpu::util::TextureBlitter,
     egui_renderer: EguiRenderer,
 
-    pipeline: RenderPipeline,
+    //pipeline: RenderPipeline,
+    pipeline: ComputePipeline,
+    surface_bg_layout: BindGroupLayout,
     voxel_cache: Arc<Mutex<VoxelCache<4>>>,
+    cache_interface: Arc<CacheProducerInterface<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>>>,
+    rm_time: Arc<Mutex<f32>>,
     insertion_debounce: bool,
 
     camera: Camera,
@@ -137,6 +120,12 @@ struct Immediates
     view_proj: Mat4,
     cam_pos: Vec3,
     frame_timestamp: u32,
+    width: u32,
+    height: u32,
+
+    chunk_width: u32,
+    chunk_height: u32,
+    chunk_alt: u32,
 }
 
 #[derive(Debug, Clone, Copy, Zeroable, Pod)]
@@ -157,6 +146,7 @@ impl State
             backends: wgpu::Backends::VULKAN,
             display: Some(Box::new(display)),
 
+            //flags: InstanceFlags::default() | InstanceFlags::debugging(),
             flags: InstanceFlags::default(),
             memory_budget_thresholds: MemoryBudgetThresholds::default(),
             backend_options: Default::default(),
@@ -167,7 +157,12 @@ impl State
             .unwrap();
         let (device, queue) = adapter
             .request_device(&wgpu::DeviceDescriptor {
-                required_features: Features::IMMEDIATES | Features::SHADER_EARLY_DEPTH_TEST,
+                required_features: Features::IMMEDIATES
+                    | Features::SHADER_EARLY_DEPTH_TEST
+                    | Features::TIMESTAMP_QUERY
+                    | Features::SHADER_I16
+                    | Features::SHADER_F16
+                    | Features::SHADER_INT64,
                 required_limits: wgpu::Limits {
                     max_immediate_size: 112,
                     max_storage_buffers_per_shader_stage: 16,
@@ -187,8 +182,21 @@ impl State
         let egui_renderer = EguiRenderer::new(&device, surface_format, &window);
 
         let mut voxel_cache = VoxelCache::<4>::new(100_000, device.clone(), queue.clone());
+        let cache_interface = CacheProducerInterface::new(256, &device);
 
         let terrain_generator = TerrainGenerator::<4>::new(5, "vxls_height.tif", 0.2, "img.jpg");
+        // let terrain_generator = TerrainGenerator::<4>::new(
+        //     5,
+        //     "./pointe_percee/height.tif",
+        //     0.2,
+        //     "./pointe_percee/ortho.jpg",
+        // );
+        // let terrain_generator = TerrainGenerator::<4>::new(
+        //     5,
+        //     "/home/albin/Documents/vxls_maps/lapiz/height.tif",
+        //     0.2,
+        //     "/home/albin/Documents/vxls_maps/lapiz/ortho.jpg",
+        // );
 
         let mut chunk_pos_map = HashMap::new();
         let chunk_instances = (0..terrain_generator.chunk_width)
@@ -212,25 +220,85 @@ impl State
         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,
+            usage: BufferUsages::COPY_DST | BufferUsages::VERTEX | BufferUsages::STORAGE,
         });
 
-        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 shader_module = unsafe {
+        //     device.create_shader_module_trusted(
+        //         wgpu::ShaderModuleDescriptor {
+        //             label: Some("Main shader module"),
+        //             source: wgpu::ShaderSource::Wgsl(
+        //                 std::fs::read_to_string("shaders/voxel.wgsl")
+        //                     .unwrap()
+        //                     .into(),
+        //             ),
+        //         },
+        //         wgpu::ShaderRuntimeChecks {
+        //             bounds_checks: false,
+        //             force_loop_bounding: false,
+        //             ray_query_initialization_tracking: false,
+        //             task_shader_dispatch_tracking: false,
+        //             mesh_shader_primitive_indices_clamp: false,
+        //             int_div_checks: false,
+        //         },
+        //     )
+        // };
+
+        //let shader_module = device.create_shader_module(include_wgsl!("../shaders/voxel.wgsl"));
+        let shader_module = unsafe {
+            device.create_shader_module_trusted(
+                wgpu::ShaderModuleDescriptor {
+                    label: Some("../shaders/voxel.spv"),
+                    source: wgpu::ShaderSource::SpirV(wgpu::__macro_helpers::Cow::Borrowed(
+                        wgpu::include_spirv_source!("../shaders/voxel.spv"),
+                    )),
+                },
+                wgpu::ShaderRuntimeChecks {
+                    bounds_checks: false,
+                    force_loop_bounding: false,
+                    ray_query_initialization_tracking: false,
+                    task_shader_dispatch_tracking: false,
+                    mesh_shader_primitive_indices_clamp: false,
+                    int_div_checks: false,
+                },
+            )
+        };
+
+        let surface_bind_group_layout =
+            device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+                label: Some("surface_bg_layout"),
+                entries: &[wgpu::BindGroupLayoutEntry {
+                    binding: 0,
+                    visibility: ShaderStages::COMPUTE,
+                    ty: wgpu::BindingType::StorageTexture {
+                        access: wgpu::StorageTextureAccess::WriteOnly,
+                        format: wgpu::TextureFormat::Rgba32Float,
+                        view_dimension: wgpu::TextureViewDimension::D2,
+                    },
+                    count: None,
+                }],
+            });
 
         let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
             label: Some("Voxel pipeline layout"),
 
-            bind_group_layouts: &[Some(&voxel_cache.bind_group_layout())],
+            bind_group_layouts: &[
+                Some(&voxel_cache.bind_group_layout()),
+                Some(&surface_bind_group_layout),
+            ],
             immediate_size: size_of::<Immediates>() as u32,
         });
 
+        let chunk_pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+            label: Some("Compute render"),
+            layout: Some(&pipeline_layout),
+            module: &shader_module,
+            entry_point: Some("ray_march_compute"),
+            compilation_options: wgpu::PipelineCompilationOptions::default(),
+            cache: None,
+        });
+
+        /*
         let chunk_pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
             label: Some("Render pipeline"),
             layout: Some(&pipeline_layout),
@@ -267,7 +335,7 @@ impl State
             depth_stencil: Some(wgpu::DepthStencilState {
                 format: wgpu::TextureFormat::Depth24PlusStencil8,
                 depth_write_enabled: Some(true),
-                depth_compare: Some(wgpu::CompareFunction::Less),
+                depth_compare: Some(wgpu::CompareFunction::LessEqual),
                 stencil: wgpu::StencilState::default(),
                 bias: wgpu::DepthBiasState::default(),
             }),
@@ -285,6 +353,7 @@ impl State
             multiview_mask: None,
             cache: None,
         });
+        */
 
         let state = State {
             instance,
@@ -294,17 +363,24 @@ impl State
             surface_format,
             egui_renderer,
             depth_buffer: Self::create_depth_buffer(&device, size.width, size.height),
+            target_texture: Self::create_target_texture(&device, size.width, size.height),
+            target_blitter: wgpu::util::TextureBlitterBuilder::new(&device, surface_format)
+                .sample_type(wgpu::FilterMode::Nearest)
+                .build(),
             queue,
             device,
+            surface_bg_layout: surface_bind_group_layout,
             usage_vec: Arc::new(Mutex::new(vec![])),
             pipeline: chunk_pipeline,
             voxel_cache: Arc::new(Mutex::new(voxel_cache)),
+            cache_interface: cache_interface.into(),
             insertion_debounce: false,
             camera: Default::default(),
             instance_buffer,
             instance_count,
             terrain_generator: Arc::new(terrain_generator),
             chunk_pos_map: chunk_pos_map.into(),
+            rm_time: Arc::new(Mutex::new(0.)),
         };
 
         // Configure surface for the first time
@@ -313,6 +389,27 @@ impl State
         state
     }
 
+    fn create_target_texture(device: &Device, width: u32, height: u32) -> (Texture, TextureView)
+    {
+        let texture = device.create_texture(&wgpu::TextureDescriptor {
+            label: Some("Target texture"),
+            size: Extent3d {
+                width,
+                height,
+                depth_or_array_layers: 1,
+            },
+            mip_level_count: 1,
+            sample_count: 1,
+            dimension: wgpu::TextureDimension::D2,
+            format: wgpu::TextureFormat::Rgba32Float,
+            usage: TextureUsages::STORAGE_BINDING | TextureUsages::TEXTURE_BINDING,
+            view_formats: &[],
+        });
+
+        let texture_view = texture.create_view(&wgpu::wgt::TextureViewDescriptor::default());
+        (texture, texture_view)
+    }
+
     fn get_window(&self) -> &Window
     {
         &self.window
@@ -388,6 +485,8 @@ impl State
         self.configure_surface();
         self.depth_buffer =
             Self::create_depth_buffer(&self.device, new_size.width, new_size.height);
+        self.target_texture =
+            Self::create_target_texture(&self.device, new_size.width, new_size.height);
     }
 
     fn render(&mut self)
@@ -398,6 +497,7 @@ impl State
         // Create texture view.
         // NOTE: We must handle Timeout because the surface may be unavailable
         // (e.g., when the window is occluded on macOS).
+        // ~~ Texture view creation ~~
         let surface_texture = match self.surface.get_current_texture()
         {
             wgpu::CurrentSurfaceTexture::Success(texture) => texture,
@@ -434,9 +534,73 @@ impl State
                 ..Default::default()
             });
 
-        // Renders a GREEN screen
-        let mut encoder = self.device.create_command_encoder(&Default::default());
+        let surface_bg = self.device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label: Some("surface_bg"),
+            layout: &self.surface_bg_layout,
+            entries: &[wgpu::BindGroupEntry {
+                binding: 0,
+                resource: wgpu::BindingResource::TextureView(&self.target_texture.1),
+            }],
+        });
 
+        // ~~ Ray-marching timestamp query setup ~~
+        let timestamp_query = self.device.create_query_set(&wgpu::QuerySetDescriptor {
+            label: Some("timestamp_query_set"),
+            ty: wgpu::QueryType::Timestamp,
+            count: 2,
+        });
+
+        let timestamp_buffer = self.device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("timestamp_buffer"),
+            size: (size_of::<u64>() * 2) as u64,
+            usage: BufferUsages::QUERY_RESOLVE | BufferUsages::COPY_SRC,
+            mapped_at_creation: false,
+        });
+
+        // ~~ Main render pass ~~
+        let mut encoder = self.device.create_command_encoder(&Default::default());
+        {
+            let mut renderpass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
+                label: Some("compute_render_pass"),
+                timestamp_writes: Some(wgpu::ComputePassTimestampWrites {
+                    query_set: &timestamp_query,
+                    beginning_of_pass_write_index: Some(0),
+                    end_of_pass_write_index: Some(1),
+                }),
+            });
+
+            renderpass.set_pipeline(&self.pipeline);
+            renderpass.set_bind_group(0, Some(&self.voxel_cache.lock().bind_group()), &[]);
+            renderpass.set_bind_group(1, Some(&surface_bg), &[]);
+            let imm = [Immediates {
+                view_proj: self.camera.view_proj().inverse(),
+                cam_pos: self.camera.position,
+                frame_timestamp: self.voxel_cache.lock().current_timestamp(),
+                width: self.size.width,
+                height: self.size.height,
+
+                chunk_width: self.terrain_generator.chunk_width as u32,
+                chunk_height: self.terrain_generator.chunk_height as u32,
+                chunk_alt: self.terrain_generator.chunk_alt as u32,
+            }];
+            renderpass.set_immediates(0, unsafe { as_raw_bytes(&imm) });
+            renderpass.dispatch_workgroups(
+                self.size.width.div_ceil(16),
+                self.size.height.div_ceil(16),
+                1,
+            );
+            drop(renderpass);
+
+            encoder.resolve_query_set(&timestamp_query, 0..2, &timestamp_buffer, 0);
+        }
+
+        self.target_blitter.copy(
+            &self.device,
+            &mut encoder,
+            &self.target_texture.1,
+            &texture_view,
+        );
+        /*
         {
             let mut renderpass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
                 label: None,
@@ -445,7 +609,12 @@ impl State
                     depth_slice: None,
                     resolve_target: None,
                     ops: wgpu::Operations {
-                        load: wgpu::LoadOp::Clear(wgpu::Color::BLACK),
+                        load: wgpu::LoadOp::Clear(wgpu::Color {
+                            r: 0.,
+                            g: 2. / 255.,
+                            b: 15. / 255.,
+                            a: 1.,
+                        }),
                         store: wgpu::StoreOp::Store,
                     },
                 })],
@@ -457,7 +626,11 @@ impl State
                     }),
                     stencil_ops: None,
                 }),
-                timestamp_writes: None,
+                timestamp_writes: Some(wgpu::RenderPassTimestampWrites {
+                    query_set: &timestamp_query,
+                    beginning_of_pass_write_index: Some(0),
+                    end_of_pass_write_index: Some(1),
+                }),
                 occlusion_query_set: None,
                 multiview_mask: None,
             });
@@ -475,26 +648,12 @@ impl State
 
             // End the renderpass.
             drop(renderpass);
+
+            encoder.resolve_query_set(&timestamp_query, 0..2, &timestamp_buffer, 0);
         }
+        */
 
-        let requests = self.device.create_buffer(&wgpu::BufferDescriptor {
-            label: Some("dummy_dumb_dinky_aaaahhh_buffer"),
-            size: 16 * 1024,
-            usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
-            mapped_at_creation: false,
-        });
-
-        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.
+        // ~~ EGUI Render pass ~~
         {
             self.egui_renderer.begin_frame(&self.window);
             egui::Window::new("Window ! ").resizable(true).show(
@@ -511,6 +670,11 @@ impl State
                                 .size(28.),
                         );
                     }
+
+                    ui.label(format!(
+                        "Ray-marching time: {}",
+                        *self.rm_time.lock() / 1_000_000.
+                    ));
                     egui_plot::Plot::new("Plot").show(ui, |plot_ui| {
                         plot_ui.bar_chart(BarChart::new(
                             "histo",
@@ -541,7 +705,7 @@ impl State
             );
         }
 
-        // Report usage
+        // ~~ Build usage buffer histogram
         let time_stamp = self.voxel_cache.lock().current_timestamp();
         let cloned_usage_histogram = self.usage_vec.clone();
         DownloadBuffer::read_buffer(
@@ -561,58 +725,75 @@ impl State
             },
         );
 
-        // Submit the command in the queue to execute
-        self.queue.submit([encoder.finish()]);
-        self.window.pre_present_notify();
-        self.queue.present(surface_texture);
+        // ~~ Do cache managment
         if !self
             .camera
             .pressed_keyset
             .contains(&winit::keyboard::KeyCode::KeyF)
-            && self.insertion_debounce
         {
-            self.insertion_debounce = false;
+            self.voxel_cache
+                .lock()
+                .cache_post_render(&mut encoder, &self.cache_interface);
         }
 
-        // if (self
-        //     .camera
-        //     .pressed_keyset
-        //     .contains(&winit::keyboard::KeyCode::KeyF))
-        //     && !self.insertion_debounce
-        // {
+        // ~~ Submit command buffer ~~
+        self.queue.submit([encoder.finish()]);
+        self.window.pre_present_notify();
+        self.queue.present(surface_texture);
+
+        // ~~ Get Ray-marching timestamps, report time ~~
+        let cloned_rm_time = self.rm_time.clone();
+        let cloned_queue = self.queue.clone();
+        DownloadBuffer::read_buffer(
+            &self.device,
+            &self.queue,
+            &timestamp_buffer.slice(..),
+            move |buffer| {
+                let buffer_slice = buffer.unwrap();
+                let slice: &[u64] = cast_slice(&buffer_slice);
+                let time = (slice[1] - slice[0]) as f32 * cloned_queue.get_timestamp_period();
+                *cloned_rm_time.lock() = time;
+            },
+        );
+
+        // ~~ Do cache managment
         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 request_count = self
+                .cache_interface
+                .total_request_count(&self.device, &self.queue);
             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 cloned_cache_interface = self.cache_interface.clone();
             let (tx, rx) = sync_channel(1);
+
+            // ~~ Download request buffer, fullfill requests, writeback to cache ~~
             DownloadBuffer::read_buffer(
                 &self.device.clone(),
                 &self.queue.clone(),
-                &requests.slice(0..),
+                &self.cache_interface.requests_buffer().slice(..),
                 move |buffer| {
                     if tx.try_send(()).is_err()
                     {
                         return;
                     }
 
-                    let cache_node_requests: Vec<CacheNodeRequest> =
+                    let cache_node_requests: Vec<CacheRequest> =
                         bytemuck::pod_collect_to_vec(&buffer.unwrap());
 
                     let generator = cloned_generator;
                     let gen_test = SineGenerator::<4>::new(5);
 
-                    let mut structure_nodes = vec![];
+                    let mut structure_nodes: Vec<StructurePoolElement<4>> = vec![];
                     let mut color_nodes: Vec<[ColorBytes; 64]> = vec![];
                     let mut location_nodes = vec![];
-                    let mut destinations = vec![];
 
                     cache_node_requests
                         .par_iter()
@@ -624,7 +805,7 @@ impl State
 
                             let location;
                             let node;
-                            if request.child_index == u32::MAX
+                            if request.locator == 0
                             {
                                 // Produce root node
                                 node =
@@ -638,9 +819,8 @@ impl State
                             else
                             {
                                 // Figure out depth of request
-                                let locator = NTreeNodeLocator::<4>::from_usize(
-                                    request.structure_locator as usize,
-                                );
+                                let locator =
+                                    NTreeNodeLocator::<4>::from_usize(request.locator as usize);
                                 let depth = locator.depth();
                                 //let (x, y, z) = locator.node_location();
 
@@ -671,13 +851,13 @@ impl State
                             }
 
                             (
-                                node.structure,
+                                StructurePoolElement {
+                                    occupancy_low: 0,
+                                    occupancy_high: 0,
+                                    pointers: node.structure,
+                                },
                                 node.colors,
                                 location,
-                                DestinationElement {
-                                    node: request.node_index,
-                                    child: request.child_index,
-                                },
                             )
                         })
                         .collect::<Vec<_>>()
@@ -686,46 +866,30 @@ impl State
                             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
                     {
-                        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,
-                            });
+                        cloned_queue.write_buffer(
+                            cloned_cache_interface.structure_nodes_buffer(),
+                            0,
+                            unsafe { as_raw_bytes(structure_nodes.as_slice()) },
+                        );
+                        cloned_queue.write_buffer(
+                            cloned_cache_interface.color_nodes_buffer(),
+                            0,
+                            unsafe { as_raw_bytes(&color_nodes.as_slice()) },
+                        );
+                        cloned_queue.write_buffer(
+                            cloned_cache_interface.location_nodes_buffer(),
+                            0,
+                            unsafe { as_raw_bytes(&&location_nodes.as_slice()) },
+                        );
 
                         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_cache
+                            .lock()
+                            .cache_insert(&mut encoder, &cloned_cache_interface);
                         cloned_queue.submit([encoder.finish()]);
                     }
                 },

src/producers.rs

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

src/sparse_tree.rs

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

src/voxel.rs

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

src/voxel/gpu.rs

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

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

src/voxel_cache.rs

@@ -0,0 +1,934 @@
+use bytemuck::Zeroable;
+use wgpu::{BindGroup, BindGroupLayout, Buffer, BufferUsages, CommandEncoder, ComputePipeline, Device, Queue, ShaderStages};
+
+use crate::voxel_cache::{data::{CacheNodeRequest, CacheResponse, ColorPoolElement, LocationPoolElement, StructurePoolElement}, producer_interface::CacheProducerInterface, request_buffer::RequestBuffer, structure_table::StructureTable, usage_buffer::UsageBuffer};
+
+
+pub mod request_buffer;
+pub mod structure_table;
+pub mod usage_buffer;
+pub mod producer_interface;
+pub mod indirect_buffer;
+
+pub mod data;
+
+pub struct VoxelCache<const N: usize>
+{
+    size: usize,
+    device: Device,
+    queue: Queue,
+    structure_pool: Buffer,
+    color_pool: Buffer,
+    location_pool: Buffer,
+
+    pub structure_table: StructureTable,
+
+    request_buffer: RequestBuffer<N>,
+
+    voxel_cache_bind_group: BindGroup,
+    structure_table_bind_group_layout: BindGroupLayout,
+    voxel_cache_render_bind_group_layout: BindGroupLayout,
+
+    // Writing requests
+    request_write_pipeline: ComputePipeline,
+
+    // User response -> caching
+    caching_pipeline: ComputePipeline,
+
+    // Invalidation
+    invalidation_pipeline: ComputePipeline,
+
+    pub usage_buffer: UsageBuffer,
+}
+
+
+#[derive(Zeroable, bytemuck::Pod, Clone, Copy)]
+#[repr(C)]
+struct CachingPipelineImmediates
+{
+    frame_timestamp: u32,
+    write_pointers: u32, // Boolean
+}
+
+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) }
+}
+
+impl<const N: usize> VoxelCache<N>
+where
+    [(); N * N * N]:,
+{
+    pub fn new(cache_size: usize, device: Device, queue: Queue) -> Self
+    {
+        let structure_pool = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("Structure pool"),
+            size: (size_of::<StructurePoolElement<N>>() * cache_size) as u64,
+            usage: BufferUsages::STORAGE,
+            mapped_at_creation: false,
+        });
+
+        let location_pool = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("Location pool"),
+            size: (size_of::<LocationPoolElement>() * cache_size) as u64,
+            usage: BufferUsages::STORAGE,
+            mapped_at_creation: false,
+        });
+
+        let color_pool = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("Color pool"),
+            size: (size_of::<ColorPoolElement<N>>() * cache_size) as u64,
+            usage: BufferUsages::STORAGE,
+            mapped_at_creation: false,
+        });
+
+        let request_buffer = RequestBuffer::new(cache_size, device.clone());
+        let usage_buffer = UsageBuffer::new(cache_size, device.clone());
+
+        let request_interface_bind_group_layout = CacheProducerInterface::<N>::request_side_bind_group_layout(&device);
+
+        // Rendering bing group layouts
+        let voxel_cache_render_bind_group_layout =
+            device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+                label: Some("voxel_cache_render_bind_group_layout"),
+                entries: &[
+                    // Structure pool
+                    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,
+                    },
+                    // Color pool
+                    wgpu::BindGroupLayoutEntry {
+                        binding: 1,
+                        visibility: ShaderStages::COMPUTE,
+                        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::COMPUTE,
+                        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::COMPUTE,
+                        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::COMPUTE,
+                        ty: wgpu::BindingType::Buffer {
+                            ty: wgpu::BufferBindingType::Storage { read_only: false },
+                            has_dynamic_offset: false,
+                            min_binding_size: None,
+                        },
+                        count: None,
+                    },
+
+                    // Structure table stuff
+                    // Structure table pointers
+                    wgpu::BindGroupLayoutEntry {
+                        binding: 5,
+                        visibility: ShaderStages::COMPUTE,
+                        ty: wgpu::BindingType::Buffer {
+                            ty: wgpu::BufferBindingType::Storage { read_only: false },
+                            has_dynamic_offset: false,
+                            min_binding_size: None,
+                        },
+                        count: None,
+                    },
+
+                    // Structure table request buffer
+                    wgpu::BindGroupLayoutEntry {
+                        binding: 6,
+                        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 voxel_cache_bind_group_layout =
+            device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+                label: Some("voxel_cache_bind_group_layout"),
+                entries: &[
+                    // Structure pool
+                    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,
+                    },
+                    // Color pool
+                    wgpu::BindGroupLayoutEntry {
+                        binding: 1,
+                        visibility: ShaderStages::COMPUTE,
+                        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::COMPUTE,
+                        ty: wgpu::BindingType::Buffer {
+                            ty: wgpu::BufferBindingType::Storage { read_only: false },
+                            has_dynamic_offset: false,
+                            min_binding_size: None,
+                        },
+                        count: None,
+                    },
+                    // Sorted requests
+                    wgpu::BindGroupLayoutEntry {
+                        binding: 3,
+                        visibility: ShaderStages::COMPUTE,
+                        ty: wgpu::BindingType::Buffer {
+                            ty: wgpu::BufferBindingType::Storage { read_only: true },
+                            has_dynamic_offset: false,
+                            min_binding_size: None,
+                        },
+                        count: None,
+                    },
+                    // LRU List
+                    wgpu::BindGroupLayoutEntry {
+                        binding: 4,
+                        visibility: ShaderStages::COMPUTE,
+                        ty: wgpu::BindingType::Buffer {
+                            ty: wgpu::BufferBindingType::Storage { read_only: true },
+                            has_dynamic_offset: false,
+                            min_binding_size: None,
+                        },
+                        count: None,
+                    },
+                    // Usage buffer
+                    wgpu::BindGroupLayoutEntry {
+                        binding: 5,
+                        visibility: ShaderStages::COMPUTE,
+                        ty: wgpu::BindingType::Buffer {
+                            ty: wgpu::BufferBindingType::Storage { read_only: false },
+                            has_dynamic_offset: false,
+                            min_binding_size: None,
+                        },
+                        count: None,
+                    },
+                    // pools request count
+                    wgpu::BindGroupLayoutEntry {
+                        binding: 6,
+                        visibility: ShaderStages::COMPUTE,
+                        ty: wgpu::BindingType::Buffer {
+                            ty: wgpu::BufferBindingType::Storage { read_only: true },
+                            has_dynamic_offset: false,
+                            min_binding_size: None,
+                        },
+                        count: None,
+                    },
+                ],
+            });
+
+        let voxel_cache_bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label: Some("voxel_cache_bind_group"),
+            layout: &voxel_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(),
+                },
+                // Sorted requests
+                wgpu::BindGroupEntry {
+                    binding: 3,
+                    resource: request_buffer.sort_buffer().as_entire_binding(),
+                },
+                // LRU list
+                wgpu::BindGroupEntry {
+                    binding: 4,
+                    resource: usage_buffer.sort_buffer().as_entire_binding(),
+                },
+                // Usage buffer
+                wgpu::BindGroupEntry {
+                    binding: 5,
+                    resource: usage_buffer.usage_buffer().as_entire_binding(),
+                },
+                // Pools request count
+                wgpu::BindGroupEntry {
+                    binding: 6,
+                    resource: request_buffer.request_count_buffer().as_entire_binding(),
+                },
+            ],
+        });
+
+        let structure_table_bind_group_layout =
+            device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+                label: Some("structure_table_bind_group_layout"),
+                entries: &[
+                    // Pointer table
+                    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,
+                    },
+                    // Sorted requests
+                    wgpu::BindGroupLayoutEntry {
+                        binding: 1,
+                        visibility: ShaderStages::COMPUTE,
+                        ty: wgpu::BindingType::Buffer {
+                            ty: wgpu::BufferBindingType::Storage { read_only: false },
+                            has_dynamic_offset: false,
+                            min_binding_size: None,
+                        },
+                        count: None,
+                    },
+                    // Requests count
+                    wgpu::BindGroupLayoutEntry {
+                        binding: 2,
+                        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 children_count = N*N*N;
+        let wgsl_bindings = 
+            format!("
+                struct StructurePoolElement
+                {{
+                     occupancy_low: atomic<u32>,
+                     occupancy_high: atomic<u32>,
+                     pointers: array<u32, {children_count}>
+                }}
+
+                struct ColorPoolElement
+                {{
+                      colors: array<u32, {children_count}>
+                }}
+
+                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> sorted_requests: array<SortedRequestsElement>;
+                @group(0) @binding(4) var<storage, read> lru_list: array<u32>;
+                @group(0) @binding(5) var<storage, read_write> usage_buffer: array<u32>;
+                @group(0) @binding(6) var<storage, read> pools_request_count: u32;
+
+                @group(1) @binding(0) var<storage, read_write> structure_table_pointers: array<u32>;
+                @group(1) @binding(1) var<storage, read_write> structure_table_sorted_requests: array<SortedRequestsElement>;
+                @group(1) @binding(2) var<storage, read_write> structure_table_request_count: u32;
+            ");
+
+        let write_requests_pipeline_layout =
+            device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+                label: Some("write_requests_pipeline_layout"),
+                bind_group_layouts: &[
+                    Some(&voxel_cache_bind_group_layout),
+                    Some(&structure_table_bind_group_layout),
+                    Some(&request_interface_bind_group_layout),
+                ],
+                immediate_size: 0,
+            });
+
+        // The user bindgroup will be created on the fly
+        let write_requests =
+            device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+                label: Some("write_requests_pipeline"),
+                layout: Some(&write_requests_pipeline_layout),
+                module: &device.create_shader_module(wgpu::ShaderModuleDescriptor {
+                    label: Some("write_requests_pipeline_shader_module"),
+                    source: wgpu::ShaderSource::Wgsl(
+                        format!("
+                            {wgsl_bindings}
+
+                            struct CacheInterfaceRequest
+                            {{
+                                structure_id: u32,
+                                locator: u32,
+                                child_index: u32
+                            }}
+
+                            struct CacheInterfaceRequestWb
+                            {{
+                                node_index: u32,
+                                child_index: u32
+                            }}
+
+                            @group(2) @binding(0) var<storage, read_write> requests: array<CacheInterfaceRequest>; 
+                            @group(2) @binding(1) var<storage, read_write> requests_wb: array<CacheInterfaceRequestWb>; 
+                            @group(2) @binding(2) var<storage, read_write> request_count: u32; 
+
+                            @group(2) @binding(3) var<storage, read> structure_nodes: array<StructurePoolElement>; 
+                            @group(2) @binding(4) var<storage, read> color_nodes: array<ColorPoolElement>; 
+                            @group(2) @binding(5) var<storage, read> locations: array<LocationPoolElement>; 
+
+                            @compute
+                            @workgroup_size(64)
+                            fn main(
+                                    @builtin(global_invocation_id) global_invocation_id: vec3<u32>
+                                )
+                            {{
+                                 // One shader invocation per invocation on both structure table domain
+                                 // and cache domain
+                                 var index = global_invocation_id.x;
+                                 let max_requests_count = min(arrayLength(&requests), arrayLength(&lru_list));
+                                 request_count = min(max_requests_count, pools_request_count + structure_table_request_count);
+                                 if(index >= request_count)
+                                 {{
+                                      return;
+                                 }}
+
+                                 if(index < structure_table_request_count)
+                                 {{
+                                     var request: CacheInterfaceRequest;
+                                     // Index in structure table IS structure id
+                                     request.structure_id = structure_table_sorted_requests[index].node;
+                                     request.locator = 0; // Root request -> locator 0
+                                     request.child_index = 0xFFFFFFFF; // child index u32::MAX ->
+
+                                     var request_wb: CacheInterfaceRequestWb;
+                                     request_wb.node_index = request.structure_id;
+                                     request_wb.child_index = 0xFFFFFFFF; // child index u32::MAX ->
+                                                                       // Structure table request
+                                     // Write request
+                                     requests[index] = request;
+                                     requests_wb[index] = request_wb;
+                                     return;
+                                 }}
+
+                                 let pool_index = index - structure_table_request_count;
+                                     
+                                 // Request location comes from location pool
+                                 var request: CacheInterfaceRequest;
+                                 request.structure_id = location_pool[sorted_requests[pool_index].node].structure_id;
+                                 request.locator = location_pool[sorted_requests[pool_index].node].structure_locator;
+                                 request.child_index = sorted_requests[pool_index].child;
+
+                                 var request_wb: CacheInterfaceRequestWb;
+                                 request_wb.node_index = sorted_requests[pool_index].node;
+                                 request_wb.child_index = sorted_requests[pool_index].child;
+
+                                 // Write request
+                                 requests[index] = request;
+                                 requests_wb[index] = request_wb;
+
+                            }}
+                        "
+                        )
+                        .into(),
+                    ),
+                }),
+                entry_point: Some("main"),
+                compilation_options: Default::default(),
+                cache: None,
+            });
+
+        // ~~~ Caching pipeline ~~~
+        // Pipeline that takes user fulling ~some~ requests
+        // and writes them to the cache based on the eviction list
+
+        let caching_pipeline_layout =
+            device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+                label: Some("caching_pipelien_layout"),
+                bind_group_layouts: &[
+                    Some(&voxel_cache_bind_group_layout),
+                    Some(&structure_table_bind_group_layout),
+                    Some(&request_interface_bind_group_layout),
+                ],
+                immediate_size: size_of::<CachingPipelineImmediates>() as u32, // Current frame timestamp
+            });
+
+        let children_count = N * N * N;
+        let caching_pipeline =
+            device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+                label: Some("caching_pipeline"),
+                layout: Some(&caching_pipeline_layout),
+                module: &device.create_shader_module(wgpu::ShaderModuleDescriptor {
+                    label: Some("caching_pipeline_shader_module"),
+                    source: wgpu::ShaderSource::Wgsl(
+                        format!("
+
+                            struct DestinationElement
+                            {{
+                                node: u32,
+                                child: u32
+                            }}
+
+                            struct CachingPipelineImmediates
+                            {{
+                                frame_timestamp: u32,
+                                write_pointers: u32
+                            }}
+
+                            {wgsl_bindings}
+                            
+                            var<immediate> parameters: CachingPipelineImmediates;
+
+                            struct CacheInterfaceRequest
+                            {{
+                                structure_id: u32,
+                                locator: u32,
+                                child_index: u32
+                            }}
+
+                            struct CacheInterfaceRequestWb
+                            {{
+                                node_index: u32,
+                                child_index: u32
+                            }}
+
+                            @group(2) @binding(0) var<storage, read_write> requests: array<CacheInterfaceRequest>; 
+                            @group(2) @binding(1) var<storage, read_write> requests_wb: array<CacheInterfaceRequestWb>; 
+                            @group(2) @binding(2) var<storage, read_write> request_count: u32; 
+
+                            @group(2) @binding(3) var<storage, read_write> structure_nodes: array<StructurePoolElement>; 
+                            @group(2) @binding(4) var<storage, read_write> color_nodes: array<ColorPoolElement>; 
+                            @group(2) @binding(5) var<storage, read_write> locations: array<LocationPoolElement>; 
+
+                            @compute
+                            @workgroup_size(64)
+                            fn main(
+                                    @builtin(global_invocation_id) global_invocation_id: vec3<u32>
+                                )
+                            {{
+                                    // Copy with indirection
+                                    let index = global_invocation_id.x;
+                                    let total = arrayLength(&structure_nodes);
+                                    let total_cache = arrayLength(&lru_list);
+                                    if(index >= total || index >= total_cache)
+                                    {{
+                                        return;
+                                    }}
+
+                                    let overwritten_element = lru_list[index];
+                                    if(parameters.write_pointers == 0 && usage_buffer[overwritten_element] != parameters.frame_timestamp)
+                                    {{
+                                        // Phase 1
+                                        // Copy into cache page
+                                         var occupancy_high = u32(0);
+                                        var occupancy_low = u32(0);
+                                        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);
+
+                                            let turn_on_bit = 
+                                                select(
+                                                u32(0),
+                                                u32(u32(1) << u32(i % ({children_count} / 2))),
+                                                ((color_nodes[index].colors[i] >> 24) != 0) ||
+                                                structure_nodes[index].pointers[i] != 0
+                                            );
+
+                                            if(i > {children_count} / 2)
+                                            {{
+                                                 occupancy_high |= turn_on_bit;
+                                            }}else
+                                            {{
+                                                 occupancy_low |= turn_on_bit;
+                                            }}
+                                        }}
+                                        atomicStore(&structure_pool[overwritten_element].occupancy_low, occupancy_low);
+                                        atomicStore(&structure_pool[overwritten_element].occupancy_high, occupancy_high);
+
+                                        color_pool[overwritten_element] = color_nodes[index];
+                                        location_pool[overwritten_element] = locations[index];
+
+                                        // Mark dirty/correct timestamp
+                                        usage_buffer[overwritten_element] = parameters.frame_timestamp + 1; 
+
+                                    }}
+
+                                    if(parameters.write_pointers != 0 && usage_buffer[overwritten_element] != parameters.frame_timestamp)
+                                    {{
+                                        // Phase 2
+                                        
+                                        // Point parent to new page
+                                        let new_pointer = (1 << 31) | (1 << 30) | overwritten_element;
+                                        if(requests_wb[index].child_index == 0xFFFFFFFF)
+                                        {{
+                                            structure_table_pointers[requests_wb[index].node_index] = new_pointer;
+                                        }}else if usage_buffer[requests_wb[index].node_index] != parameters.frame_timestamp + 1
+                                        {{
+                                            structure_pool[requests_wb[index].node_index].pointers[requests_wb[index].child_index] = new_pointer;
+
+                                            let turn_on_bit = u32(1 << (requests_wb[index].child_index % ({children_count} / 2)));
+                                            if(requests_wb[index].child_index > {children_count} / 2)
+                                            {{
+                                                //atomicOr(&structure_pool[requests_wb[index].node_index].occupancy_high, turn_on_bit);
+                                            }}else
+                                            {{
+                                                //atomicOr(&structure_pool[requests_wb[index].node_index].occupancy_low, turn_on_bit);
+                                            }}
+                                        }}
+                                    }}
+
+                            }}
+                        ")
+                        .into(),
+                    ),
+                }),
+                entry_point: Some("main"),
+                compilation_options: Default::default(),
+                cache: None,
+            });
+
+        // ~~~ Invalidation pipeline ~~~
+        // Pipeline invalidates old nodes that points to ones
+        // that have been replaces base on timestamp information
+        let invalidation_pipeline_layout =
+            device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+                label: Some("invalidation_pipeline_layout"),
+                bind_group_layouts: &[
+                    Some(&voxel_cache_bind_group_layout),
+                    Some(&structure_table_bind_group_layout),
+                ],
+                immediate_size: size_of::<u32>() as u32, // frame_timestamp
+            });
+
+        let invalidation_pipeline =
+            device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+                label: Some("invalidation_pipeline"),
+                layout: Some(&invalidation_pipeline_layout),
+                module: &device.create_shader_module(wgpu::ShaderModuleDescriptor {
+                    label: Some("invalidation_pipeline_shader_module"),
+                    source: wgpu::ShaderSource::Wgsl(
+                        format!("
+                            {wgsl_bindings}
+                            
+                            var<immediate> frame_timestamp: u32;
+
+                            @compute
+                            @workgroup_size(64)
+                            fn main(
+                                    @builtin(global_invocation_id) global_invocation_id: vec3<u32>
+                                )
+                            {{
+                                    var index = global_invocation_id.x;
+                                    let total_pool = arrayLength(&structure_pool);
+                                    let total_table = arrayLength(&structure_table_pointers);
+
+                                    if index < total_pool
+                                    {{
+                                         for(var i = 0; i < {children_count}; i += 1)
+                                         {{
+                                            let structure_pointer = structure_pool[index].pointers[i];
+                                            let pointer = structure_pointer & 0x3FFFFFFF;
+                                            let pointer_subdiv = ((structure_pointer >> 31) & 1) != 0;
+                                            let pointed_timestamp = usage_buffer[pointer];
+                                            if(pointed_timestamp == frame_timestamp + 1) // Future
+                                                                                         // timestamp
+                                                                                         // -> new page
+                                            {{
+                                                 // Invalidate
+                                                 structure_pool[index].pointers[i] = select(u32(0), u32(1), pointer_subdiv) << 31;
+                                            }}
+                                         }}
+                                         return;
+                                    }} 
+                                    index -= total_pool;
+                                    if index < total_table
+                                    {{ 
+                                            let structure_pointer = structure_table_pointers[index];
+                                            // Is pointer pointing to something valid
+                                            let pointer_valid = ((structure_pointer >> 30) & 1) != 0;
+                                            let pointer_subdiv = ((structure_pointer >> 31) & 1) != 0;
+                                            let pointer = structure_pointer & 0x3FFFFFFF;
+                                            let pointed_timestamp = usage_buffer[pointer];
+                                            if(pointed_timestamp == frame_timestamp + 1 && pointer_valid && pointer_subdiv) // Future
+                                                                                         // timestamp
+                                                                                         // -> new page
+                                            {{
+                                                 // Invalidate
+                                                 structure_table_pointers[index] &= select(u32(0), u32(1), pointer_subdiv) << 31;
+                                            }}
+                                    }}
+
+                            }}
+                        ")
+                        .into(),
+                    ),
+                }),
+                entry_point: Some("main"),
+                compilation_options: Default::default(),
+                cache: None,
+            });
+
+        Self {
+            size: cache_size,
+            structure_table: StructureTable::new(device.clone(), queue.clone()),
+            queue,
+            device,
+            structure_pool,
+            color_pool,
+            location_pool,
+            request_buffer,
+            usage_buffer,
+
+            voxel_cache_bind_group,
+            structure_table_bind_group_layout,
+            voxel_cache_render_bind_group_layout,
+
+            // Write requests stage
+            request_write_pipeline: write_requests,
+
+            // Caching operation
+            caching_pipeline,
+
+            // Invalidation
+            invalidation_pipeline,
+        }
+    }
+
+    pub fn bind_group_layout(&self) -> BindGroupLayout
+    {
+        self.voxel_cache_render_bind_group_layout.clone()
+    }
+
+    pub fn bind_group(&self) -> BindGroup
+    {
+        self.device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label: Some("voxel_cache_render_bind_group"),
+            layout: &self.voxel_cache_render_bind_group_layout,
+            entries: &[
+                // Structure pool
+                wgpu::BindGroupEntry {
+                    binding: 0,
+                    resource: self.structure_pool.as_entire_binding(),
+                },
+                // Color pool
+                wgpu::BindGroupEntry {
+                    binding: 1,
+                    resource: self.color_pool.as_entire_binding(),
+                },
+                // Location pool
+                wgpu::BindGroupEntry {
+                    binding: 2,
+                    resource: self.location_pool.as_entire_binding(),
+                },
+                // Request buffer 
+                wgpu::BindGroupEntry {
+                    binding: 3,
+                    resource: self.request_buffer.request_buffer().as_entire_binding(),
+                },
+                // Usage buffer
+                wgpu::BindGroupEntry {
+                    binding: 4,
+                    resource: self.usage_buffer.usage_buffer().as_entire_binding(),
+                },
+                // Structure table
+                // Structure table pointers
+                wgpu::BindGroupEntry {
+                    binding: 5,
+                    resource: self.structure_table.pointer_table.as_entire_binding(),
+                },
+                // Pools request count
+                wgpu::BindGroupEntry {
+                    binding: 6,
+                    resource: self.structure_table.request_buffer.request_buffer().as_entire_binding(),
+                },
+            ],
+        })
+    }
+
+    fn structure_table_bind_group(&self) -> BindGroup
+    {
+        self.device.create_bind_group(&wgpu::BindGroupDescriptor
+            {
+                label: Some("structure_table_bind_group"),
+                layout: &self.structure_table_bind_group_layout,
+                entries: 
+                &[
+                    wgpu::BindGroupEntry
+                    {
+                        binding: 0,
+                        resource: self.structure_table.pointer_table.as_entire_binding(),
+                    },
+                    wgpu::BindGroupEntry
+                    {
+                        binding: 1,
+                        resource: self.structure_table.request_buffer.sort_buffer().as_entire_binding(),
+                    },
+                    wgpu::BindGroupEntry
+                    {
+                        binding: 2,
+                        resource: self.structure_table.request_buffer.request_count_buffer().as_entire_binding(),
+                    },
+                ],
+            })
+    }
+
+    pub fn cache_post_render(&mut self, encoder: &mut CommandEncoder, request_interface: &CacheProducerInterface<N>)
+    {
+        // Sort requests, reset request buffers to count requests
+        self.usage_buffer.sort_usage(encoder);
+        self.request_buffer.sort_requests(encoder);
+
+        self.structure_table
+            .request_buffer_mut()
+            .sort_requests(encoder);
+
+        // Sorted requests are now in the group
+        let mut write_requests_pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
+            label: Some("write_requests_pass"),
+            timestamp_writes: None,
+        });
+
+        write_requests_pass.set_bind_group(0, Some(&self.voxel_cache_bind_group), &[]);
+        write_requests_pass.set_bind_group(1, Some(&self.structure_table_bind_group()), &[]);
+        write_requests_pass.set_bind_group(2, Some(&request_interface.request_side_bind_group), &[]);
+        write_requests_pass.set_pipeline(&self.request_write_pipeline);
+
+        // Compute necessary shader invocations
+        let shader_invocation_count = request_interface.size;
+        let workgroup_invocation_count = shader_invocation_count.div_ceil(64);
+        write_requests_pass.dispatch_workgroups(workgroup_invocation_count as u32, 1, 1);
+        drop(write_requests_pass);
+
+        self.request_buffer.reset_requests(encoder);
+        self.structure_table
+            .request_buffer_mut()
+            .reset_requests(encoder);
+    }
+
+    pub fn current_timestamp(&self) -> u32
+    {
+        self.usage_buffer.timestamp()
+    }
+
+    pub fn cache_insert(&mut self, encoder: &mut CommandEncoder, cache_interface: &CacheProducerInterface<N>)
+    {
+        // Each of the entry of the cache insertion, is matched with the usage list to evict
+        let mut cache_insertion_pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
+            label: Some("cache_insertion_pass"),
+            timestamp_writes: None,
+        });
+
+        // Copy insertions to evicted lru
+        // Point parents to children
+        let structure_table_bind_group = self.structure_table_bind_group();
+
+        // ~~~ Write data + Dirty flagging/timestamp update ~~~
+        {
+            cache_insertion_pass.set_pipeline(&self.caching_pipeline);
+            cache_insertion_pass.set_bind_group(0, Some(&self.voxel_cache_bind_group), &[]);
+            cache_insertion_pass.set_bind_group(1, Some(&structure_table_bind_group), &[]);
+            cache_insertion_pass.set_bind_group(2, Some(&cache_interface.request_side_bind_group), &[]);
+            cache_insertion_pass.set_immediates(
+                0,
+                bytemuck::bytes_of(&CachingPipelineImmediates {
+                    frame_timestamp: self.usage_buffer.timestamp(),
+                    write_pointers: 0, // false
+                }),
+            );
+
+            // Compute dispatch amounts
+            let shader_invocation_count = cache_interface.size;
+            let workgroup_invocations = shader_invocation_count.div_ceil(64);
+
+            cache_insertion_pass.dispatch_workgroups(workgroup_invocations as u32, 1, 1);
+        }
+
+        // ~~~ Invalidate pointers ~~~
+        {
+            cache_insertion_pass.set_pipeline(&self.invalidation_pipeline);
+            cache_insertion_pass.set_bind_group(2, None, &[]);
+            cache_insertion_pass
+                .set_immediates(0, bytemuck::bytes_of(&self.usage_buffer.timestamp()));
+
+            // Compute dispatch amounts
+            let shader_invocation_count = self.size + self.structure_table.allocation_table.len();
+            let workgroup_invocations = shader_invocation_count.div_ceil(64);
+
+            cache_insertion_pass.dispatch_workgroups(workgroup_invocations as u32, 1, 1);
+        }
+
+        // ~~~ Write new pointers ~~~
+        {
+            cache_insertion_pass.set_pipeline(&self.caching_pipeline);
+            cache_insertion_pass.set_bind_group(2, Some(&cache_interface.request_side_bind_group), &[]);
+            cache_insertion_pass.set_immediates(
+                0,
+                bytemuck::bytes_of(&CachingPipelineImmediates {
+                    frame_timestamp: self.usage_buffer.timestamp(),
+                    write_pointers: 1, // true
+                }),
+            );
+
+            // Compute dispatch amounts
+            let shader_invocation_count = cache_interface.size;
+            let workgroup_invocations = shader_invocation_count.div_ceil(64);
+
+            cache_insertion_pass.dispatch_workgroups(workgroup_invocations as u32, 1, 1);
+        }
+
+    }
+
+    pub fn next_frame(&mut self)
+    {
+        self.usage_buffer.next_frame();
+    }
+
+}

src/voxel_cache/data.rs

@@ -0,0 +1,105 @@
+use bytemuck::Pod;
+use bytemuck::Zeroable;
+use wgpu::Buffer;
+
+use crate::sparse_tree::Color;
+
+#[derive(Clone, Copy, Zeroable, Pod, Debug)]
+#[repr(C)]
+pub struct CacheNodeRequest
+{
+    // Requested ressource
+    pub structure_id: u32,
+    pub structure_locator: u32,
+
+    // Write back info
+    pub node_index: u32,
+    pub child_index: u32,
+}
+
+pub struct RequestBufferElement<const N: usize>
+where
+    [(); N * N * N]:,
+{
+    request_count: [u32; N * N * N],
+}
+
+#[repr(C)]
+pub struct StructurePoolElement<const N: usize>
+where
+    [(); N * N * N]:,
+{
+    pub occupancy_low: u32,
+    pub occupancy_high: u32,
+    pub pointers: [StructurePointer; N * N * N],
+}
+
+pub struct DestinationElement
+{
+    pub node: u32,
+    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: [ColorBytes; N * N * N],
+}
+
+pub struct LocationPoolElement
+{
+    pub structure_id: u32,
+    pub structure_locator: u32,
+}
+
+pub struct CacheResponse
+{
+    // Each buffer contains the same amount of elements (structure of arrays style)
+
+    // Cache data to bring in
+    pub structure_nodes: Buffer,
+    pub color_nodes: Buffer,
+    pub locations: Buffer,
+
+    // Which nodes this extends : node_index + child_index
+    pub parents: Buffer,
+}
+
+#[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_cache/indirect_buffer.rs

@@ -0,0 +1,44 @@
+use wgpu::{Buffer, BufferUsages, Device, util::DeviceExt};
+
+pub struct BufferCompactor
+{
+    size: usize,
+    block_count: usize,
+    reduced_buffer: Buffer,
+    sum_buffer: Buffer,
+    compaction_buffer: Buffer,
+}
+
+impl BufferCompactor
+{
+    const THREAD_COUNT: usize = 256;
+    pub fn new(device: &Device, size: usize) -> Self
+    {
+        let block_count = size.div_ceil(size);
+        let reduced_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
+            label: Some("indirect_buffer_reduced"),
+            contents: bytemuck::cast_slice(vec![0; block_count].as_slice()),
+            usage: BufferUsages::STORAGE,
+        });
+
+        let sum_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
+            label: Some("indirect_buffer_sum"),
+            contents: bytemuck::cast_slice(vec![0; size].as_slice()),
+            usage: BufferUsages::STORAGE,
+        });
+
+        let compaction_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
+            label: Some("indirect_buffer_compaction"),
+            contents: bytemuck::cast_slice(vec![0; size].as_slice()),
+            usage: BufferUsages::STORAGE,
+        });
+
+        BufferCompactor {
+            size,
+            block_count,
+            reduced_buffer,
+            sum_buffer,
+            compaction_buffer,
+        }
+    }
+}

src/voxel_cache/producer_interface.rs

@@ -0,0 +1,369 @@
+// The cache emits requests, and receives answer to these requests
+// The cache producer interface provides the necessary buffers to contains
+// these, so that use applications can fullfill the requests
+
+use bytemuck::Pod;
+use bytemuck::Zeroable;
+use wgpu::BindGroup;
+use wgpu::BindGroupEntry;
+use wgpu::BindGroupLayout;
+use wgpu::BindGroupLayoutEntry;
+use wgpu::Buffer;
+use wgpu::BufferUsages;
+use wgpu::CommandEncoder;
+use wgpu::Device;
+use wgpu::Queue;
+use wgpu::util::DownloadBuffer;
+
+use crate::voxel_cache::data::ColorPoolElement;
+use crate::voxel_cache::data::LocationPoolElement;
+use crate::voxel_cache::data::StructurePoolElement;
+
+#[derive(Clone, Copy, Zeroable, Pod)]
+#[repr(C)]
+pub struct CacheRequest
+{
+    pub structure_id: u32,
+    pub locator: u32,
+    pub child_index: u32,
+}
+
+pub(super) struct CacheRequestWbInfo
+{
+    node_index: u32,
+    child_index: u32,
+}
+
+pub struct CacheProducerInterface<const N: usize>
+where
+    [(); N * N * N]:,
+{
+    // Request emissions
+    pub(super) size: usize,
+
+    // Contains the specific node being requested
+    pub(super) requests: Buffer,
+
+    // Contains how many request have effectively been emitted
+    pub(super) request_count: Buffer,
+
+    // Contains on which node the request was done
+    pub(super) requests_wb_info: Buffer,
+    // Request fullfillment
+    pub(super) structure_nodes: Buffer,
+    pub(super) color_nodes: Buffer,
+    pub(super) location_nodes: Buffer,
+
+    pub(super) request_side_bind_group_layout: BindGroupLayout,
+    pub(super) request_side_bind_group: BindGroup,
+
+    pub(super) producer_side_bind_group_layout: BindGroupLayout,
+    pub(super) producer_side_bind_group: BindGroup,
+}
+impl<const N: usize> CacheProducerInterface<N>
+where
+    [(); N * N * N]:,
+{
+    pub fn new(size: usize, device: &Device) -> Self
+    {
+        let requests = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("cache_interface_request_buffer"),
+            size: (size_of::<CacheRequest>() * size) as u64,
+            usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
+            mapped_at_creation: false,
+        });
+
+        let requests_wb_info = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("cache_interface_writeback_info"),
+            size: (size_of::<CacheRequestWbInfo>() * size) as u64,
+            usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
+            mapped_at_creation: false,
+        });
+
+        let structure_nodes = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("cache_interface_structure_nodes"),
+            size: (size_of::<StructurePoolElement<N>>() * size) as u64,
+            usage: BufferUsages::STORAGE | BufferUsages::COPY_DST,
+            mapped_at_creation: false,
+        });
+
+        let color_nodes = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("cache_interface_structure_nodes"),
+            size: (size_of::<ColorPoolElement<N>>() * size) as u64,
+            usage: BufferUsages::STORAGE | BufferUsages::COPY_DST,
+            mapped_at_creation: false,
+        });
+
+        let location_nodes = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("cache_interface_structure_nodes"),
+            size: (size_of::<LocationPoolElement>() * size) as u64,
+            usage: BufferUsages::STORAGE | BufferUsages::COPY_DST,
+            mapped_at_creation: false,
+        });
+
+        let request_count = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("cache_interface_request_count"),
+            size: (size_of::<u32>()) as u64,
+            usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
+            mapped_at_creation: false,
+        });
+
+        let request_side_bind_group_layout = Self::request_side_bind_group_layout(device);
+        let producer_side_bind_group_layout =
+            device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+                label: Some("cache_interface_producer_side_bind_group_layout"),
+                entries: &[
+                    // Reqests
+                    BindGroupLayoutEntry {
+                        binding: 0,
+                        visibility: wgpu::ShaderStages::COMPUTE,
+                        ty: wgpu::BindingType::Buffer {
+                            ty: wgpu::BufferBindingType::Storage { read_only: true },
+                            has_dynamic_offset: false,
+                            min_binding_size: None,
+                        },
+                        count: None,
+                    },
+                    // Request count
+                    BindGroupLayoutEntry {
+                        binding: 1,
+                        visibility: wgpu::ShaderStages::COMPUTE,
+                        ty: wgpu::BindingType::Buffer {
+                            ty: wgpu::BufferBindingType::Storage { read_only: true },
+                            has_dynamic_offset: false,
+                            min_binding_size: None,
+                        },
+                        count: None,
+                    },
+                    // Structure nodes
+                    BindGroupLayoutEntry {
+                        binding: 2,
+                        visibility: wgpu::ShaderStages::COMPUTE,
+                        ty: wgpu::BindingType::Buffer {
+                            ty: wgpu::BufferBindingType::Storage { read_only: false },
+                            has_dynamic_offset: false,
+                            min_binding_size: None,
+                        },
+                        count: None,
+                    },
+                    // Color nodes
+                    BindGroupLayoutEntry {
+                        binding: 3,
+                        visibility: wgpu::ShaderStages::COMPUTE,
+                        ty: wgpu::BindingType::Buffer {
+                            ty: wgpu::BufferBindingType::Storage { read_only: false },
+                            has_dynamic_offset: false,
+                            min_binding_size: None,
+                        },
+                        count: None,
+                    },
+                    // Location nodes
+                    BindGroupLayoutEntry {
+                        binding: 4,
+                        visibility: wgpu::ShaderStages::COMPUTE,
+                        ty: wgpu::BindingType::Buffer {
+                            ty: wgpu::BufferBindingType::Storage { read_only: false },
+                            has_dynamic_offset: false,
+                            min_binding_size: None,
+                        },
+                        count: None,
+                    },
+                ],
+            });
+
+        // Bind groups
+        let request_side_bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label: Some("cache_interface_request_side_bind_group"),
+            layout: &request_side_bind_group_layout,
+            entries: &[
+                BindGroupEntry {
+                    binding: 0,
+                    resource: requests.as_entire_binding(),
+                },
+                BindGroupEntry {
+                    binding: 1,
+                    resource: requests_wb_info.as_entire_binding(),
+                },
+                BindGroupEntry {
+                    binding: 2,
+                    resource: request_count.as_entire_binding(),
+                },
+                BindGroupEntry {
+                    binding: 3,
+                    resource: structure_nodes.as_entire_binding(),
+                },
+                BindGroupEntry {
+                    binding: 4,
+                    resource: color_nodes.as_entire_binding(),
+                },
+                BindGroupEntry {
+                    binding: 5,
+                    resource: location_nodes.as_entire_binding(),
+                },
+            ],
+        });
+
+        let producer_side_bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label: Some("cache_interface_producer_side_bind_group"),
+            layout: &producer_side_bind_group_layout,
+            entries: &[
+                BindGroupEntry {
+                    binding: 0,
+                    resource: requests.as_entire_binding(),
+                },
+                BindGroupEntry {
+                    binding: 1,
+                    resource: request_count.as_entire_binding(),
+                },
+                BindGroupEntry {
+                    binding: 2,
+                    resource: structure_nodes.as_entire_binding(),
+                },
+                BindGroupEntry {
+                    binding: 3,
+                    resource: color_nodes.as_entire_binding(),
+                },
+                BindGroupEntry {
+                    binding: 4,
+                    resource: location_nodes.as_entire_binding(),
+                },
+            ],
+        });
+
+        Self {
+            size,
+            requests,
+            request_count,
+            requests_wb_info,
+            structure_nodes,
+            color_nodes,
+            location_nodes,
+
+            request_side_bind_group_layout,
+            request_side_bind_group,
+            producer_side_bind_group_layout,
+            producer_side_bind_group,
+        }
+    }
+
+    pub fn requests_buffer(&self) -> &Buffer
+    {
+        &self.requests
+    }
+
+    pub fn structure_nodes_buffer(&self) -> &Buffer
+    {
+        &self.structure_nodes
+    }
+
+    pub fn color_nodes_buffer(&self) -> &Buffer
+    {
+        &self.color_nodes
+    }
+
+    pub fn location_nodes_buffer(&self) -> &Buffer
+    {
+        &self.location_nodes
+    }
+
+    pub fn total_request_count(&self, device: &Device, queue: &Queue) -> u32
+    {
+        let (tx, rx) = std::sync::mpsc::sync_channel(1);
+        wgpu::util::DownloadBuffer::read_buffer(
+            device,
+            queue,
+            &self.request_count.slice(0..),
+            move |download_buffer| {
+                let vec = download_buffer.unwrap().to_vec();
+                let value = u32::from_ne_bytes(std::array::from_fn(|i| vec[i]));
+                let _ = tx.try_send(value);
+            },
+        );
+
+        loop
+        {
+            if let Ok(value) = rx.try_recv()
+            {
+                return value;
+            }
+            device.poll(wgpu::wgt::PollType::Poll).unwrap();
+        }
+
+        //panic!("Could not retrieve total request count.");
+    }
+
+    pub fn request_side_bind_group_layout(device: &Device) -> BindGroupLayout
+    {
+        device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+            label: Some("cache_interface_request_side_bind_group_layout"),
+            entries: &[
+                // Reqests
+                BindGroupLayoutEntry {
+                    binding: 0,
+                    visibility: wgpu::ShaderStages::COMPUTE,
+                    ty: wgpu::BindingType::Buffer {
+                        ty: wgpu::BufferBindingType::Storage { read_only: false },
+                        has_dynamic_offset: false,
+                        min_binding_size: None,
+                    },
+                    count: None,
+                },
+                // Requests writeback
+                BindGroupLayoutEntry {
+                    binding: 1,
+                    visibility: wgpu::ShaderStages::COMPUTE,
+                    ty: wgpu::BindingType::Buffer {
+                        ty: wgpu::BufferBindingType::Storage { read_only: false },
+                        has_dynamic_offset: false,
+                        min_binding_size: None,
+                    },
+                    count: None,
+                },
+                // Request count
+                BindGroupLayoutEntry {
+                    binding: 2,
+                    visibility: wgpu::ShaderStages::COMPUTE,
+                    ty: wgpu::BindingType::Buffer {
+                        ty: wgpu::BufferBindingType::Storage { read_only: false },
+                        has_dynamic_offset: false,
+                        min_binding_size: None,
+                    },
+                    count: None,
+                },
+                // Structure nodes
+                BindGroupLayoutEntry {
+                    binding: 3,
+                    visibility: wgpu::ShaderStages::COMPUTE,
+                    ty: wgpu::BindingType::Buffer {
+                        ty: wgpu::BufferBindingType::Storage { read_only: false },
+                        has_dynamic_offset: false,
+                        min_binding_size: None,
+                    },
+                    count: None,
+                },
+                // Color nodes
+                BindGroupLayoutEntry {
+                    binding: 4,
+                    visibility: wgpu::ShaderStages::COMPUTE,
+                    ty: wgpu::BindingType::Buffer {
+                        ty: wgpu::BufferBindingType::Storage { read_only: false },
+                        has_dynamic_offset: false,
+                        min_binding_size: None,
+                    },
+                    count: None,
+                },
+                // Location nodes
+                BindGroupLayoutEntry {
+                    binding: 5,
+                    visibility: wgpu::ShaderStages::COMPUTE,
+                    ty: wgpu::BindingType::Buffer {
+                        ty: wgpu::BufferBindingType::Storage { read_only: false },
+                        has_dynamic_offset: false,
+                        min_binding_size: None,
+                    },
+                    count: None,
+                },
+            ],
+        })
+    }
+}

src/voxel_cache/request_buffer.rs

@@ -0,0 +1,641 @@
+use bytemuck::bytes_of;
+use wgpu::BindGroup;
+use wgpu::BindGroupEntry;
+use wgpu::Buffer;
+use wgpu::BufferUsages;
+use wgpu::CommandEncoder;
+use wgpu::ComputePipeline;
+use wgpu::Device;
+use wgpu::ShaderStages;
+
+use crate::voxel_cache::data::RequestBufferElement;
+
+pub struct RequestBuffer<const N: usize>
+{
+    pub cache_size: usize,
+    pub request_buffer: Buffer,
+    pub request_count_buffer: Buffer,
+    pub indirect_count_storage: Buffer,
+    pub element_sort_count_buffer: Buffer,
+    pub indirect_count: Buffer,
+    pub sort_buffer: Buffer,
+    pub reset_pipeline: ComputePipeline,
+    pub sort_pipeline: ComputePipeline,
+    pub bindgroup: BindGroup,
+    pub device: Device,
+
+    pub compaction_pipeline: ComputePipeline,
+    pub running_sum_pipeline: ComputePipeline,
+}
+
+impl<const N: usize> RequestBuffer<N>
+where
+    [(); N * N * N]:,
+{
+    pub fn new(cache_size: usize, device: Device) -> Self
+    {
+        let request_buffer = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some(format!("request_buffer_{N}").as_str()),
+            size: (size_of::<RequestBufferElement<N>>() * cache_size) as u64,
+            usage: BufferUsages::STORAGE | BufferUsages::COPY_DST,
+            mapped_at_creation: false,
+        });
+
+        let request_count_buffer = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("Request_count_buffer"),
+            size: size_of::<u32>() as u64,
+            usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
+            mapped_at_creation: false,
+        });
+
+        let indirect_count = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("indirect_count_buffer"),
+            usage: BufferUsages::INDIRECT | BufferUsages::COPY_DST,
+            size: size_of::<wgpu::util::DispatchIndirectArgs>() as u64,
+            mapped_at_creation: false,
+        });
+
+        let indirect_count_storage = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("indirect_count_storage_buffer"),
+            usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
+            size: size_of::<wgpu::util::DispatchIndirectArgs>() as u64,
+            mapped_at_creation: false,
+        });
+
+        let element_sort_count_buffer = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("indirect_count_storage_buffer"),
+            usage: BufferUsages::STORAGE,
+            size: size_of::<u32>() as u64,
+            mapped_at_creation: false,
+        });
+
+        // One element number per child entry
+        // One number to identify node, one to identify sub child
+        let sort_buffer_count = cache_size * N * N * N;
+        let sort_buffer = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some(format!("request_sort_buffer_{N}").as_str()),
+            size: (size_of::<(u32, u32)>() * sort_buffer_count) as u64,
+            usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
+            mapped_at_creation: true,
+        });
+
+        let init_data = (0..(cache_size as u32))
+            .flat_map(|i| (0..((N * N * N) as u32)).map(move |j| (i, j)))
+            .flat_map(|(i, j)| [i, j]) // Because bytemuck does not like tuples ...
+            .collect::<Vec<_>>();
+
+        sort_buffer
+            .get_mapped_range_mut(0..)
+            .unwrap()
+            .copy_from_slice(bytemuck::cast_slice(&init_data));
+        sort_buffer.unmap();
+
+        let bindgroup_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+            label: Some("request_sort_pipeline_bing_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,
+                },
+                wgpu::BindGroupLayoutEntry {
+                    binding: 1,
+                    visibility: ShaderStages::COMPUTE,
+                    ty: wgpu::BindingType::Buffer {
+                        ty: wgpu::BufferBindingType::Storage { read_only: false },
+                        has_dynamic_offset: false,
+                        min_binding_size: None,
+                    },
+                    count: None,
+                },
+                wgpu::BindGroupLayoutEntry {
+                    binding: 2,
+                    visibility: ShaderStages::COMPUTE,
+                    ty: wgpu::BindingType::Buffer {
+                        ty: wgpu::BufferBindingType::Storage { read_only: false },
+                        has_dynamic_offset: false,
+                        min_binding_size: None,
+                    },
+                    count: None,
+                },
+                wgpu::BindGroupLayoutEntry {
+                    binding: 3,
+                    visibility: ShaderStages::COMPUTE,
+                    ty: wgpu::BindingType::Buffer {
+                        ty: wgpu::BufferBindingType::Storage { read_only: false },
+                        has_dynamic_offset: false,
+                        min_binding_size: None,
+                    },
+                    count: None,
+                },
+                wgpu::BindGroupLayoutEntry {
+                    binding: 4,
+                    visibility: ShaderStages::COMPUTE,
+                    ty: wgpu::BindingType::Buffer {
+                        ty: wgpu::BufferBindingType::Storage { read_only: false },
+                        has_dynamic_offset: false,
+                        min_binding_size: None,
+                    },
+                    count: None,
+                },
+            ],
+        });
+
+        let bindgroup = device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label: Some("request_sort_pipeline_bindgroup"),
+            layout: &bindgroup_layout,
+            entries: &[
+                BindGroupEntry {
+                    binding: 0,
+                    resource: request_buffer.as_entire_binding(),
+                },
+                BindGroupEntry {
+                    binding: 1,
+                    resource: sort_buffer.as_entire_binding(),
+                },
+                BindGroupEntry {
+                    binding: 2,
+                    resource: request_count_buffer.as_entire_binding(),
+                },
+                BindGroupEntry {
+                    binding: 3,
+                    resource: indirect_count_storage.as_entire_binding(),
+                },
+                BindGroupEntry {
+                    binding: 4,
+                    resource: element_sort_count_buffer.as_entire_binding(),
+                },
+            ],
+        });
+
+        let pipeline_layouts = &device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+            label: Some("request_buffer_pipeline_layout"),
+            bind_group_layouts: &[Some(&bindgroup_layout)],
+            immediate_size: size_of::<u32>() as u32,
+        });
+
+        let children_count = N * N * N;
+        let sort_pipeline =
+            device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+                label: Some("Sort node request"),
+                layout: Some(pipeline_layouts),
+                module: &device.create_shader_module(wgpu::ShaderModuleDescriptor {
+                    label: Some("clear_chunk_requests shader module"),
+                    source: wgpu::ShaderSource::Wgsl(
+                        format!("
+                            struct RequestElement
+                            {{
+                                children: array<u32, {children_count}> 
+                            }}
+
+                            struct SortElement
+                            {{
+                                 node: u32,
+                                 child: u32
+                            }}
+
+                            var<immediate> both_phase: u32;
+
+                            @group(0) @binding(0) var<storage, read_write> request_buffer: array<RequestElement>;
+                            @group(0) @binding(1) var<storage, read_write> sort_indirection: array<SortElement>; 
+                            @group(0) @binding(2) var<storage, read_write> request_count: u32; 
+                            @group(0) @binding(3) var<storage, read_write> indirect_count: vec3<u32>;
+                            @group(0) @binding(4) var<storage, read_write> element_sort_count: u32; 
+
+                            fn big_fusion(index: u32, phase_size: u32) -> vec2<u32>
+                            {{
+                                // Find out in which block index this invocation pertains
+                                let block_index = index / (phase_size / 2);
+                                let element_index = index % (phase_size / 2); 
+
+                                let offset = block_index * phase_size;
+                                let a = offset + element_index;
+                                let b = offset + phase_size - 1 - element_index;
+                                return vec2<u32>(a, b);
+                            }}
+
+                            fn small_fusion(index: u32, phase: u32, sub_phase: u32) -> vec2<u32>
+                            {{
+                                let phase_size = u32(1 << (phase - sub_phase + 1));
+                                
+                                let block_index = index / (phase_size / 2);
+                                let element_index = index % (phase_size / 2); 
+
+                                let offset = block_index * phase_size;
+                                let a = offset + element_index;
+                                let b = a + (phase_size / 2);
+                                return vec2<u32>(a, b);
+                            }}
+
+
+                            @compute
+                            @workgroup_size(64)
+                            fn main(
+                                    @builtin(global_invocation_id) global_invocation_id: vec3<u32>
+                                )
+                            {{
+
+                                 let index = global_invocation_id.x;
+                                 let total = arrayLength(&sort_indirection);
+
+                                 let sub_phase = (both_phase >> 16) & 0xFFFF;
+                                 let phase = both_phase & 0xFFFF;
+                                 let phase_total_width = u32(1 << (phase + 1));
+
+
+                                 // Check if phase is last
+                                 if(phase == sub_phase)
+                                 {{
+                                      let next_phase_width = phase_total_width * 2;
+                                      if(next_phase_width >= (element_sort_count * 2) && phase_total_width >= (element_sort_count * 2))
+                                      {{
+                                        indirect_count = vec3<u32>(0);
+                                      }}
+                                 }}
+
+                                 if(phase == 0)
+                                 {{
+                                    let a = index * 2;
+                                    let b = a + 1;
+                                    sort_indirection[a].child = a % {children_count};
+                                    sort_indirection[b].child = b % {children_count};
+                                 }}
+
+                                 var swap_indices = vec2<u32>(0, 0);
+
+                                 if(sub_phase == 0)
+                                 {{
+                                    // Bitonic fusion
+                                    swap_indices = big_fusion(index, phase_total_width);
+                                 }}else
+                                 {{
+                                    swap_indices = small_fusion(index, phase, sub_phase);
+                                 }}
+
+                                 // Do swap
+
+                                 //if(swap_indices.y >= request_count_round_up)
+                                 if(swap_indices.y >= element_sort_count)
+                                 {{
+                                    // Suppose that swap_indices.y is -inf, dont swap
+                                    return;
+                                 }}
+                                 
+                                 let av = request_buffer[sort_indirection[swap_indices.x].node].children[sort_indirection[swap_indices.x].child];
+                                 let bv = request_buffer[sort_indirection[swap_indices.y].node].children[sort_indirection[swap_indices.y].child];
+
+                                 if(bv > av)
+                                 {{
+                                    let temp = sort_indirection[swap_indices.x];
+                                    sort_indirection[swap_indices.x] = sort_indirection[swap_indices.y];
+                                    sort_indirection[swap_indices.y] = temp;
+                                 }}
+                            }}
+                        ")
+                        .into(),
+                    ),
+                }),
+                entry_point: Some("main"),
+                compilation_options: Default::default(),
+                cache: None,
+            });
+
+        let running_sum_pipeline =
+            device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+                label: Some("running_sum_pipeline"),
+                layout: Some(pipeline_layouts),
+                module: &device.create_shader_module(wgpu::ShaderModuleDescriptor {
+                    label: Some("running_sum shader module"),
+                    source: wgpu::ShaderSource::Wgsl(
+                        format!("
+                            struct RequestElement
+                            {{
+                                children: array<u32, {children_count}> 
+                            }}
+
+                            struct SortElement
+                            {{
+                                 node: u32,
+                                 child: u32
+                            }}
+
+                            var<immediate> phase: u32;
+
+                            @group(0) @binding(0) var<storage, read_write> request_buffer: array<RequestElement>;
+                            @group(0) @binding(1) var<storage, read_write> sort_indirection: array<SortElement>; 
+                            @group(0) @binding(2) var<storage, read_write> request_counts: atomic<u32>; 
+                            @group(0) @binding(3) var<storage, read_write> indirect_count: u32; 
+                            @group(0) @binding(4) var<storage, read_write> node_sort_count: u32; 
+
+                            @compute
+                            @workgroup_size(64)
+                            fn main(
+                                    @builtin(global_invocation_id) global_invocation_id: vec3<u32>
+                                )
+                            {{
+                                let index = global_invocation_id.x;
+                                let len = arrayLength(&request_buffer);
+                                if(index > len) {{ return; }}
+                                let sindex = index * {children_count};
+                                 
+                                if(phase == 0)
+                                {{
+                                    var count = 0;
+                                    for(var i = 0; i < {children_count}; i++)
+                                    {{
+                                        count += select(0, 1, request_buffer[index].children[i] != 0);
+                                    }}
+
+                                    sort_indirection[sindex].child = select(u32(0), u32(1), count != 0);
+                                    if count != 0
+                                    {{
+                                        atomicAdd(&request_counts, 1);
+                                    }}
+                                    return;
+                                }}
+
+                                if(phase == 0xFFFFFFFF)
+                                {{
+                                    // Double buffering bring back
+                                    sort_indirection[sindex].child = sort_indirection[sindex].node;
+                                }}
+
+                                // Phase is not zero, running sum part
+                                let running_sum_phase = phase - 1;
+                                let running_sum_offset = u32((1 << running_sum_phase) * {children_count});
+
+                                var add = u32(0);
+                                // Double buffering
+                                if(running_sum_phase % 2 == 0)
+                                {{
+                                     if(sindex >= running_sum_offset)
+                                     {{ 
+                                        add = sort_indirection[sindex - running_sum_offset].child;
+                                     }}
+                                     sort_indirection[sindex].node = sort_indirection[sindex].child + add;
+                                }}
+                                else
+                                {{
+                                     if(sindex >= running_sum_offset)
+                                     {{ 
+                                        add = sort_indirection[sindex - running_sum_offset].node;
+                                     }}
+                                     sort_indirection[sindex].child = sort_indirection[sindex].node + add;
+                                }}
+                            }}
+                        ")
+                        .into(),
+                    ),
+                }),
+                entry_point: Some("main"),
+                compilation_options: Default::default(),
+                cache: None,
+            });
+
+        let compaction_pipeline =
+            device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+                label: Some("compaction_pipeline"),
+                layout: Some(pipeline_layouts),
+                module: &device.create_shader_module(wgpu::ShaderModuleDescriptor {
+                    label: Some("compaction shader module"),
+                    source: wgpu::ShaderSource::Wgsl(
+                        format!("
+                            struct RequestElement
+                            {{
+                                children: array<u32, {children_count}> 
+                            }}
+
+                            struct SortElement
+                            {{
+                                 node: u32,
+                                 child: u32
+                            }}
+
+                            var<immediate> phase: u32;
+
+                            @group(0) @binding(0) var<storage, read_write> request_buffer: array<RequestElement>;
+                            @group(0) @binding(1) var<storage, read_write> sort_indirection: array<SortElement>; 
+                            @group(0) @binding(2) var<storage, read_write> request_counts: atomic<u32>; 
+                            @group(0) @binding(3) var<storage, read_write> indirect_count: vec3<u32>;
+                            @group(0) @binding(4) var<storage, read_write> element_sort_count: u32; 
+
+                            @compute
+                            @workgroup_size(64)
+                            fn main(
+                                    @builtin(global_invocation_id) global_invocation_id: vec3<u32>
+                                )
+                            {{
+                                 let index = global_invocation_id.x;
+                                 let sindex = index * {children_count};
+                                 let len = arrayLength(&request_buffer);
+
+                                 element_sort_count = sort_indirection[(len - 1) * {children_count}].child * {children_count};
+                                 let invocation_count = (element_sort_count / 2) + select(u32(0), u32(1), element_sort_count % 2 != 0);
+                                 indirect_count = vec3(
+                                     (invocation_count / 64) + select(u32(0), u32(1), invocation_count % 64 != 0),
+                                     1, 1
+                                 );
+                                 let destination_index = sort_indirection[sindex].child - 1;
+
+                                 var count = u32(0);
+                                 for(var i = 0; i < {children_count}; i++)
+                                 {{
+                                      count += request_buffer[index].children[i];
+                                 }}
+
+                                 if(count != 0) // Keep ?
+                                 {{
+                                    for(var i = u32(0); i < u32({children_count}); i++)
+                                    {{
+                                        sort_indirection[destination_index * u32({children_count}) + i].node = index;
+                                    }}
+                                 }}
+                            }}
+                        ")
+                        .into(),
+                    ),
+                }),
+                entry_point: Some("main"),
+                compilation_options: Default::default(),
+                cache: None,
+            });
+
+        let reset_pipeline =
+            device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+                label: Some("reset_node_requests"),
+                layout: Some(pipeline_layouts),
+                module: &device.create_shader_module(wgpu::ShaderModuleDescriptor {
+                    label: Some("clear_chunk_requests shader module"),
+                    source: wgpu::ShaderSource::Wgsl(
+                        format!("
+                            struct RequestElement
+                            {{
+                                children: array<u32, {children_count}> 
+                            }}
+
+                            @group(0) @binding(0) var<storage, read_write> request_buffer: array<RequestElement>;
+                            @group(0) @binding(1) var<storage, read_write> _ignore: array<u32>; 
+                            @group(0) @binding(2) var<storage, read_write> request_count: u32;
+                            @group(0) @binding(4) var<storage, read_write> node_sort_count: u32; 
+
+                            @compute
+                            @workgroup_size(64)
+                            fn main(
+                                    @builtin(global_invocation_id) global_invocation_id: vec3<u32>
+                                )
+                            {{
+                                    let index = global_invocation_id.x;
+                                    let total = arrayLength(&request_buffer);
+                                    if(index >= total)
+                                    {{ return; }}
+
+                                    if(index < total)
+                                    {{
+                                        for(var i = 0; i < {children_count}; i += 1)
+                                        {{
+                                             /*
+                                            if(request_buffer[index].children[i] != 0)
+                                            {{
+                                                atomicAdd(&request_count, 1);
+                                            }}
+                                            */
+                                            request_buffer[index].children[i] = 0;
+                                        }}
+                                        request_count = 0;
+                                    }}
+                            }}
+                        ")
+                        .into(),
+                    ),
+                }),
+                entry_point: Some("main"),
+                compilation_options: Default::default(),
+                cache: None,
+            });
+
+        Self {
+            cache_size,
+            request_buffer,
+            request_count_buffer,
+            indirect_count,
+            indirect_count_storage,
+            element_sort_count_buffer,
+            running_sum_pipeline,
+            sort_buffer,
+            sort_pipeline,
+            reset_pipeline,
+            compaction_pipeline,
+            bindgroup,
+            device,
+        }
+    }
+
+    pub fn request_buffer(&self) -> &Buffer
+    {
+        &self.request_buffer
+    }
+
+    pub fn request_count_buffer(&self) -> &Buffer
+    {
+        &self.request_count_buffer
+    }
+
+    pub fn sort_buffer(&self) -> &Buffer
+    {
+        &self.sort_buffer
+    }
+
+    pub fn reset_requests(&self, encoder: &mut CommandEncoder)
+    {
+        let mut compute_pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
+            label: Some("request_buffer_reset_compute_pass"),
+            timestamp_writes: None,
+        });
+
+        compute_pass.set_bind_group(0, Some(&self.bindgroup), &[]);
+        compute_pass.set_pipeline(&self.reset_pipeline);
+
+        let shader_invocations = self.cache_size; // one invocation per element
+        let workgroup_invocations = shader_invocations.div_ceil(64);
+        compute_pass.dispatch_workgroups(workgroup_invocations as u32, 1, 1);
+    }
+
+    pub fn sort_requests(&self, encoder: &mut CommandEncoder)
+    {
+        let mut compaction_compute_pass =
+            encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
+                label: Some("request_buffer_compaction_compute_pass"),
+                timestamp_writes: None,
+            });
+
+        compaction_compute_pass.set_bind_group(0, Some(&self.bindgroup), &[]);
+        let request_element_count = self.cache_size; // Each child slot is sorted
+        let workgroups_invocations = request_element_count.div_ceil(64);
+
+        // = Perform list compaction
+
+        // == Running sum
+        compaction_compute_pass.set_pipeline(&self.running_sum_pipeline);
+
+        // Phase 0: Put ones in correct location
+        compaction_compute_pass.set_immediates(0, bytes_of(&0));
+        compaction_compute_pass.dispatch_workgroups(workgroups_invocations as u32, 1, 1);
+        // Phase _: running sum
+        // Running sum phase
+        let running_sum_steps = request_element_count.next_power_of_two().ilog2();
+        for i in 1..=running_sum_steps
+        {
+            compaction_compute_pass.set_immediates(0, bytes_of(&i));
+            compaction_compute_pass.dispatch_workgroups(workgroups_invocations as u32, 1, 1);
+        }
+
+        if !running_sum_steps.is_multiple_of(2)
+        {
+            // Bring back double buffer
+            compaction_compute_pass.set_immediates(0, bytes_of(&0xFFFFFFFF_u32));
+            compaction_compute_pass.dispatch_workgroups(workgroups_invocations as u32, 1, 1);
+        }
+
+        // == Stream compaction
+        compaction_compute_pass.set_pipeline(&self.compaction_pipeline);
+        compaction_compute_pass.dispatch_workgroups(workgroups_invocations as u32, 1, 1);
+        drop(compaction_compute_pass);
+
+        // == Copy count into indirect buffer
+
+        // = Sort
+
+        // == Bitonic sort
+        let sort_element_count = self.cache_size * N * N * N;
+        let phases_upper_bound = sort_element_count.next_power_of_two().ilog2();
+
+        // Phase 0 dispatch all
+        for i in 0..=phases_upper_bound
+        {
+            encoder.copy_buffer_to_buffer(
+                &self.indirect_count_storage,
+                0,
+                &self.indirect_count,
+                0,
+                Some(size_of::<wgpu::util::DispatchIndirectArgs>() as u64),
+            );
+            let mut sort_compute_pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
+                label: Some(format!("request_buffer_sort_compute_pass_{}", 0).as_str()),
+                timestamp_writes: None,
+            });
+            sort_compute_pass.set_bind_group(0, Some(&self.bindgroup), &[]);
+            sort_compute_pass.set_pipeline(&self.sort_pipeline);
+            for j in 0..=i
+            {
+                sort_compute_pass.set_immediates(0, bytemuck::bytes_of(&(i | (j << 16))));
+                sort_compute_pass.dispatch_workgroups_indirect(&self.indirect_count, 0);
+            }
+            drop(sort_compute_pass);
+        }
+    }
+}

src/voxel_cache/structure_table.rs

@@ -0,0 +1,129 @@
+use std::num::NonZero;
+
+use wgpu::Buffer;
+use wgpu::BufferUsages;
+use wgpu::CommandEncoder;
+use wgpu::Device;
+use wgpu::Queue;
+use wgpu::util::StagingBelt;
+
+use crate::voxel_cache::data::StructurePointer;
+use crate::voxel_cache::request_buffer::RequestBuffer;
+
+// Stores root pointers to the cache
+pub struct StructureTable
+{
+    pub(crate) device: Device,
+    pub(crate) queue: Queue,
+    pub(crate) allocation_table: Vec<bool>,
+    pub(crate) available_slots: usize,
+    pub(crate) pointer_table: Buffer,
+    pub(crate) request_buffer: RequestBuffer<1>,
+    pub(crate) write_staging: StagingBelt,
+}
+
+impl StructureTable
+{
+    pub fn new(device: Device, queue: Queue) -> Self
+    {
+        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 | BufferUsages::COPY_SRC,
+            mapped_at_creation: false,
+        });
+
+        StructureTable {
+            request_buffer: RequestBuffer::new(1, device.clone()),
+            write_staging: StagingBelt::new(device.clone(), size_of::<u32>() as u64),
+            device,
+            queue,
+            allocation_table: vec![false],
+            available_slots: 1,
+            pointer_table,
+        }
+    }
+
+    fn double_capacity(&mut self, encoder: &mut CommandEncoder)
+    {
+        self.available_slots += self.allocation_table.len();
+        self.allocation_table
+            .extend(vec![false; self.allocation_table.len()]);
+
+        // Copy buffers into bigger buffers
+        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 | BufferUsages::COPY_SRC | BufferUsages::COPY_DST,
+            mapped_at_creation: false,
+        });
+
+        encoder.copy_buffer_to_buffer(
+            &self.pointer_table,
+            0,
+            &pointer_table,
+            0,
+            self.pointer_table.size(),
+        );
+
+        self.pointer_table = pointer_table;
+        self.request_buffer = RequestBuffer::new(self.allocation_table.len(), self.device.clone());
+    }
+
+    pub fn request_buffer(&self) -> &RequestBuffer<1>
+    {
+        &self.request_buffer
+    }
+
+    pub fn request_buffer_mut(&mut self) -> &mut RequestBuffer<1>
+    {
+        &mut self.request_buffer
+    }
+
+    pub fn remove_structure(&mut self, structure_id: u32)
+    {
+        self.available_slots += 1;
+        self.allocation_table[structure_id as usize] = false;
+    }
+
+    // Returns a structure Id
+    pub fn allocate_structure(&mut self, subdivided: bool) -> u32
+    {
+        let mut encoder = self
+            .device
+            .create_command_encoder(&wgpu::CommandEncoderDescriptor {
+                label: Some("structure_table_resize_encoder"),
+            });
+
+        // Get first available element
+        if self.available_slots == 0
+        {
+            self.double_capacity(&mut encoder);
+        }
+        let (first_id, _) = self
+            .allocation_table
+            .iter()
+            .enumerate()
+            .filter(|(_, allocated)| !**allocated)
+            .next()
+            .unwrap();
+        self.allocation_table[first_id] = true;
+        self.available_slots -= 1;
+
+        // Write empty pointer to new pointer
+        self.write_staging
+            .write_buffer(
+                &mut encoder,
+                &self.pointer_table,
+                size_of::<u32>() as u64 * first_id as u64,
+                NonZero::new(size_of::<u32>() as u64).unwrap(),
+            )
+            .copy_from_slice(bytemuck::bytes_of(
+                &StructurePointer::new(subdivided, false, 0).0,
+            ));
+        self.write_staging.finish_and_recall_on_submit(&encoder);
+        self.queue.submit([encoder.finish()]);
+
+        first_id as u32
+    }
+}

src/voxel_cache/usage_buffer.rs

@@ -0,0 +1,249 @@
+use bytemuck::bytes_of;
+use wgpu::BindGroup;
+use wgpu::BindGroupEntry;
+use wgpu::Buffer;
+use wgpu::BufferUsages;
+use wgpu::CommandEncoder;
+use wgpu::ComputePipeline;
+use wgpu::Device;
+use wgpu::ShaderStages;
+
+pub struct UsageBuffer
+{
+    pub cache_size: usize,
+    pub current_timestamp: u32,
+    pub usage_buffer: Buffer,
+    pub sort_buffer: Buffer,
+    pub sort_pipeline: ComputePipeline,
+    pub bindgroup: BindGroup,
+    pub device: Device,
+}
+
+impl UsageBuffer
+{
+    pub fn new(cache_size: usize, device: Device) -> Self
+    {
+        let initial_timestamp = 0;
+        let usage_buffer = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("Usage buffer"),
+            size: (size_of::<u32>() * cache_size) as u64,
+            usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC,
+            mapped_at_creation: true,
+        });
+        let sort_buffer = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("usage_sort_buffer"),
+            size: (size_of::<u32>() * cache_size) as u64,
+            usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
+            mapped_at_creation: true,
+        });
+
+        usage_buffer
+            .get_mapped_range_mut(0..)
+            .unwrap()
+            .copy_from_slice(bytemuck::cast_slice(
+                vec![initial_timestamp; cache_size].as_slice(),
+            ));
+        usage_buffer.unmap();
+
+        sort_buffer
+            .get_mapped_range_mut(0..)
+            .unwrap()
+            .copy_from_slice(bytemuck::cast_slice(
+                &(0..(cache_size as u32)).collect::<Vec<_>>(),
+            ));
+        sort_buffer.unmap();
+
+        let bindgroup_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+            label: Some("usage_sort_pipeline_bing_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,
+                },
+                wgpu::BindGroupLayoutEntry {
+                    binding: 1,
+                    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 bindgroup = device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label: Some("usage_sort_pipeline_bindgroup"),
+            layout: &bindgroup_layout,
+            entries: &[
+                BindGroupEntry {
+                    binding: 0,
+                    resource: usage_buffer.as_entire_binding(),
+                },
+                BindGroupEntry {
+                    binding: 1,
+                    resource: sort_buffer.as_entire_binding(),
+                },
+            ],
+        });
+
+        let pipeline_layouts = &device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+            label: Some("usage_buffer_pipeline_layout"),
+            bind_group_layouts: &[Some(&bindgroup_layout)],
+            immediate_size: size_of::<u32>() as u32,
+        });
+
+        let sort_pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+            label: Some("Sort node request"),
+            layout: Some(pipeline_layouts),
+            module: &device.create_shader_module(wgpu::ShaderModuleDescriptor {
+                label: Some("clear_chunk_requests shader module"),
+                source: wgpu::ShaderSource::Wgsl(
+                    "
+
+                            @group(0) @binding(0) var<storage, read_write> usage_buffer: array<u32>;
+                            @group(0) @binding(1) var<storage, read_write> sort_buffer: array<u32>; 
+                            var<immediate> both_phase: u32;
+
+                            fn big_fusion(index: u32, phase_size: u32) -> vec2<u32>
+                            {{
+                                // Find out in which block index this invocation pertains
+                                let block_index = index / (phase_size / 2);
+                                let element_index = index % (phase_size / 2); 
+
+                                let offset = block_index * phase_size;
+                                let a = offset + element_index;
+                                let b = offset + phase_size - 1 - element_index;
+                                return vec2<u32>(a, b);
+                            }}
+
+                            fn small_fusion(index: u32, phase: u32, sub_phase: u32) -> vec2<u32>
+                            {{
+                                let phase_size = u32(1 << (phase - sub_phase + 1));
+                                
+                                let block_index = index / (phase_size / 2);
+                                let element_index = index % (phase_size / 2); 
+
+                                let offset = block_index * phase_size;
+                                let a = offset + element_index;
+                                let b = a + (phase_size / 2);
+                                return vec2<u32>(a, b);
+                            }}
+
+
+                            @compute
+                            @workgroup_size(64)
+                            fn main(
+                                    @builtin(global_invocation_id) global_invocation_id: vec3<u32>
+                                )
+                            {{
+
+                                 let length = arrayLength(&sort_buffer);
+                                 let index = global_invocation_id.x;
+                                 let sub_phase = (both_phase >> 16) & 0xFFFF;
+                                 let phase = both_phase & 0xFFFF;
+                                 let phase_total_width = u32(1 << (phase + 1));
+
+                                 var swap_indices = vec2<u32>(0, 0);
+
+                                 if(sub_phase == 0)
+                                 {{
+                                    // Bitonic fusion
+                                    swap_indices = big_fusion(index, phase_total_width);
+                                 }}else
+                                 {{
+                                    swap_indices = small_fusion(index, phase, sub_phase);
+                                 }}
+
+                                 // Do swap
+                                 if(swap_indices.y >= length)
+                                 {{
+                                    // Suppose that swap_indices.y is -inf, dont swap
+                                    return;
+                                 }}
+                                 
+                                 let av = usage_buffer[sort_buffer[swap_indices.x]];
+                                 let bv = usage_buffer[sort_buffer[swap_indices.y]];
+
+                                 if(bv < av)
+                                 {{
+                                    let temp = sort_buffer[swap_indices.x];
+                                    sort_buffer[swap_indices.x] = sort_buffer[swap_indices.y];
+                                    sort_buffer[swap_indices.y] = temp;
+                                 }}
+                            }}
+                        "
+                    .into(),
+                ),
+            }),
+            entry_point: Some("main"),
+            compilation_options: Default::default(),
+            cache: None,
+        });
+
+        Self {
+            cache_size,
+            current_timestamp: initial_timestamp,
+            usage_buffer,
+            sort_buffer,
+            sort_pipeline,
+            bindgroup,
+            device,
+        }
+    }
+
+    pub fn timestamp(&self) -> u32
+    {
+        self.current_timestamp
+    }
+
+    pub fn next_frame(&mut self)
+    {
+        let (new_timestamp, _) = self.current_timestamp.overflowing_add(1);
+        self.current_timestamp = new_timestamp;
+    }
+
+    pub fn usage_buffer(&self) -> &Buffer
+    {
+        &self.usage_buffer
+    }
+
+    pub fn sort_buffer(&self) -> &Buffer
+    {
+        &self.sort_buffer
+    }
+
+    pub fn sort_usage(&self, encoder: &mut CommandEncoder)
+    {
+        let mut compute_pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
+            label: Some("usage_buffer_sort_compute_pass"),
+            timestamp_writes: None,
+        });
+
+        compute_pass.set_bind_group(0, Some(&self.bindgroup), &[]);
+        compute_pass.set_pipeline(&self.sort_pipeline);
+
+        // Compute required shader invocations
+        let sort_steps = self.cache_size.next_power_of_two().ilog2(); // Each element is sorted
+        let shader_invocations = self.cache_size.div_ceil(2); // bitonic sorting :
+        // half as many shaders
+        // per element
+        let workgroup_invocations = shader_invocations.div_ceil(64);
+
+        for i in 0..=sort_steps
+        {
+            for j in 0..=i
+            {
+                compute_pass.set_immediates(0, bytes_of(&(j << 16 | i)));
+                compute_pass.dispatch_workgroups(workgroup_invocations as u32, 1, 1);
+            }
+        }
+    }
+}