diff --git a/shaders/compaction.slang b/shaders/compaction.slang
index 35828f2..3ed1b15 100644
--- a/shaders/compaction.slang
+++ b/shaders/compaction.slang
@@ -113,6 +113,7 @@ void scatter(
     uint32_t3 local_thread_id: SV_GroupThreadID,
     uint32_t3 global_thread_id: SV_DispatchThreadID)
 {
+    let total = index_buffer.getCount();
     let thread_index = global_thread_id.x;
     let local_thread_index = local_thread_id.x;
 
@@ -123,12 +124,12 @@ void scatter(
     let sum2 = sum_buffer[thread_index * 2 + 1] + reduced_value;
 
     // Scatter index to compact index buffer if predicate is true
-    if (count_buffer[thread_index * 2] != 0)
+    if (count_buffer[thread_index * 2] != 0 && sum1 < total)
     {
         index_buffer[sum1] = thread_index * 2;
     }
 
-    if (count_buffer[thread_index * 2 + 1] != 0)
+    if (count_buffer[thread_index * 2 + 1] != 0 && sum2 < total)
     {
         index_buffer[sum2] = thread_index * 2 + 1;
     }
diff --git a/shaders/compaction.spv b/shaders/compaction.spv
index 7e9b2a0..e31434a 100644
Binary files a/shaders/compaction.spv and b/shaders/compaction.spv differ
diff --git a/src/main.rs b/src/main.rs
index 73bef09..a1d9e58 100644
--- a/src/main.rs
+++ b/src/main.rs
@@ -69,7 +69,7 @@ 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::indirect_buffer::BufferCompactor;
+use crate::voxel_cache::indirect_buffer::NonZeroBufferCompactor;
 use crate::voxel_cache::producer_interface::CacheProducerInterface;
 use crate::voxel_cache::producer_interface::CacheRequest;
 
@@ -105,6 +105,7 @@ struct State
     instance_count: usize,
     usage_vec: Arc<Mutex<Vec<usize>>>,
     rm_time: Arc<Mutex<f32>>,
+    frame_time: Arc<Mutex<f32>>,
     insertion_debounce: bool,
 
     camera: Camera,
@@ -164,6 +165,7 @@ impl State
                     | Features::TIMESTAMP_QUERY
                     | Features::SHADER_I16
                     | Features::SHADER_F16
+                    | Features::TIMESTAMP_QUERY_INSIDE_ENCODERS
                     | Features::SHADER_INT64,
                 required_limits: wgpu::Limits {
                     max_immediate_size: 112,
@@ -176,94 +178,6 @@ impl State
             .unwrap();
 
         let mut encoder = device.create_command_encoder(&Default::default());
-        // --------------------------
-        let mut count_buffer = vec![0u32; 65536];
-        let mut check_buffer = vec![];
-        // Spread random numbers throughout the buffer
-        let mut picked = HashSet::new();
-        for i in 1..=1000
-        {
-            let mut r = rand::random_range(0..(count_buffer.len()));
-            while picked.contains(&r)
-            {
-                r = rand::random_range(0..(count_buffer.len()));
-            }
-
-            picked.insert(r);
-            check_buffer.push(r);
-            count_buffer[r] = i;
-        }
-
-        check_buffer.sort();
-
-        count_buffer
-            .iter()
-            .enumerate()
-            .filter(|(_, x)| **x != 0)
-            .take(100)
-            .for_each(|(x, y)| println!("{}, {}", x, y));
-
-        let buffer_compactor = BufferCompactor::new(&device, count_buffer.len());
-        let gpu_count_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
-            label: None,
-            contents: bytemuck::cast_slice(count_buffer.as_slice()),
-            usage: BufferUsages::STORAGE,
-        });
-        let gpu_index_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
-            label: None,
-            contents: bytemuck::cast_slice(vec![0u32; count_buffer.len()].as_slice()),
-            usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
-        });
-
-        buffer_compactor.compact_buffer(
-            &device,
-            &mut encoder,
-            &gpu_count_buffer,
-            &gpu_index_buffer,
-        );
-
-        queue.submit([encoder.finish()]);
-        // --------------------------
-
-        wgpu::util::DownloadBuffer::read_buffer(
-            &device,
-            &queue,
-            &buffer_compactor.reduced_buffer.slice(..),
-            //&buffer_compactor.total_sum_buffer(),
-            |result| {
-                let result = result.unwrap();
-                let indices: Vec<u32> = result
-                    .to_vec()
-                    .chunks(4)
-                    .map(|x| u32::from_ne_bytes([x[0], x[1], x[2], x[3]]))
-                    .collect();
-
-                println!("Total sum {:?}", indices);
-            },
-        );
-
-        wgpu::util::DownloadBuffer::read_buffer(
-            &device,
-            &queue,
-            &gpu_index_buffer.slice(..),
-            |result| {
-                let result = result.unwrap();
-                let indices: Vec<u32> = result
-                    .to_vec()
-                    .chunks(4)
-                    .map(|x| u32::from_ne_bytes([x[0], x[1], x[2], x[3]]))
-                    .collect();
-
-                println!("{:?}", &indices[..1010]);
-            },
-        );
-        for _ in 0..16
-        {
-            device.poll(wgpu::wgt::PollType::Poll);
-        }
-
-        println!("Check buffer : {:?}", check_buffer);
-        panic!();
 
         let size = window.inner_size();
 
@@ -473,6 +387,7 @@ impl State
             terrain_generator: Arc::new(terrain_generator),
             chunk_pos_map: chunk_pos_map.into(),
             rm_time: Arc::new(Mutex::new(0.)),
+            frame_time: Arc::new(Mutex::new(0.)),
         };
 
         // Configure surface for the first time
@@ -639,25 +554,26 @@ impl State
         let timestamp_query = self.device.create_query_set(&wgpu::QuerySetDescriptor {
             label: Some("timestamp_query_set"),
             ty: wgpu::QueryType::Timestamp,
-            count: 2,
+            count: 4,
         });
 
         let timestamp_buffer = self.device.create_buffer(&wgpu::BufferDescriptor {
             label: Some("timestamp_buffer"),
-            size: (size_of::<u64>() * 2) as u64,
+            size: (size_of::<u64>() * 4) 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());
+        encoder.write_timestamp(&timestamp_query, 0);
         {
             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),
+                    beginning_of_pass_write_index: Some(2),
+                    end_of_pass_write_index: Some(3),
                 }),
             });
 
@@ -682,8 +598,6 @@ impl State
                 1,
             );
             drop(renderpass);
-
-            encoder.resolve_query_set(&timestamp_query, 0..2, &timestamp_buffer, 0);
         }
 
         self.target_blitter.copy(
@@ -763,6 +677,10 @@ impl State
                         );
                     }
 
+                    ui.label(format!(
+                        "frame        time: {}",
+                        *self.frame_time.lock() / 1_000_000.
+                    ));
                     ui.label(format!(
                         "Ray-marching time: {}",
                         *self.rm_time.lock() / 1_000_000.
@@ -829,12 +747,15 @@ impl State
         }
 
         // ~~ Submit command buffer ~~
+        encoder.write_timestamp(&timestamp_query, 1);
+        encoder.resolve_query_set(&timestamp_query, 0..4, &timestamp_buffer, 0);
         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_frame_time = self.frame_time.clone();
         let cloned_queue = self.queue.clone();
         DownloadBuffer::read_buffer(
             &self.device,
@@ -843,8 +764,10 @@ impl State
             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;
+                let frame_time = (slice[1] - slice[0]) as f32 * cloned_queue.get_timestamp_period();
+                let rm_time = (slice[3] - slice[2]) as f32 * cloned_queue.get_timestamp_period();
+                *cloned_rm_time.lock() = rm_time;
+                *cloned_frame_time.lock() = frame_time;
             },
         );
 
diff --git a/src/voxel_cache.rs b/src/voxel_cache.rs
index 70e42a7..5c3ae0b 100644
--- a/src/voxel_cache.rs
+++ b/src/voxel_cache.rs
@@ -282,7 +282,7 @@ where
                 // Sorted requests
                 wgpu::BindGroupEntry {
                     binding: 3,
-                    resource: request_buffer.sort_buffer().as_entire_binding(),
+                    resource: request_buffer.index_buffer().as_entire_binding(),
                 },
                 // LRU list
                 wgpu::BindGroupEntry {
@@ -372,13 +372,13 @@ where
                 @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(3) var<storage, read> sorted_requests: array<u32>;
                 @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(1) var<storage, read_write> structure_table_sorted_requests: array<u32>;
                 @group(1) @binding(2) var<storage, read_write> structure_table_request_count: u32;
             ");
 
@@ -445,7 +445,7 @@ where
                                  {{
                                      var request: CacheInterfaceRequest;
                                      // Index in structure table IS structure id
-                                     request.structure_id = structure_table_sorted_requests[index].node;
+                                     request.structure_id = structure_table_sorted_requests[index];
                                      request.locator = 0; // Root request -> locator 0
                                      request.child_index = 0xFFFFFFFF; // child index u32::MAX ->
 
@@ -460,16 +460,19 @@ where
                                  }}
 
                                  let pool_index = index - structure_table_request_count;
-                                     
+                                
+                                 let sorted_requests_node = sorted_requests[pool_index] / 64;
+                                 let sorted_requests_child = sorted_requests[pool_index] % 64;
+                                 
                                  // 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;
+                                 request.structure_id = location_pool[sorted_requests_node].structure_id;
+                                 request.locator = location_pool[sorted_requests_node].structure_locator;
+                                 request.child_index = sorted_requests_child;
 
                                  var request_wb: CacheInterfaceRequestWb;
-                                 request_wb.node_index = sorted_requests[pool_index].node;
-                                 request_wb.child_index = sorted_requests[pool_index].child;
+                                 request_wb.node_index = sorted_requests_node;
+                                 request_wb.child_index = sorted_requests_child;
 
                                  // Write request
                                  requests[index] = request;
@@ -809,7 +812,7 @@ where
                     wgpu::BindGroupEntry
                     {
                         binding: 1,
-                        resource: self.structure_table.request_buffer.sort_buffer().as_entire_binding(),
+                        resource: self.structure_table.request_buffer.index_buffer().as_entire_binding(),
                     },
                     wgpu::BindGroupEntry
                     {
diff --git a/src/voxel_cache/request_buffer.rs b/src/voxel_cache/request_buffer.rs
index 78dbc4a..7f88cb3 100644
--- a/src/voxel_cache/request_buffer.rs
+++ b/src/voxel_cache/request_buffer.rs
@@ -9,23 +9,27 @@ use wgpu::Device;
 use wgpu::ShaderStages;
 
 use crate::voxel_cache::data::RequestBufferElement;
+use crate::voxel_cache::indirect_buffer::NonZeroBufferCompactor;
 
 pub struct RequestBuffer<const N: usize>
 {
+    pub device: Device,
     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 index_buffer: Buffer,
+
+    pub indirect_count_storage: Buffer,
+    pub indirect_count: Buffer,
+    pub request_count: Buffer,
+
+    pub buffer_compactor: NonZeroBufferCompactor,
+    pub bindgroup: BindGroup,
+    pub compute_indirect_pipeline: ComputePipeline,
     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>
@@ -41,13 +45,6 @@ where
             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,
@@ -62,33 +59,29 @@ where
             mapped_at_creation: false,
         });
 
-        let element_sort_count_buffer = device.create_buffer(&wgpu::BufferDescriptor {
-            label: Some("indirect_count_storage_buffer"),
-            usage: BufferUsages::STORAGE,
+        let request_count = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("request_count_storage_buffer"),
+            usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC | BufferUsages::COPY_DST,
             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,
+        let index_buffer_count = cache_size * N * N * N;
+        let index_buffer = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some(format!("request_index_buffer_{N}").as_str()),
+            size: (size_of::<u32>() * index_buffer_count) as u64,
             usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
-            mapped_at_creation: true,
+            mapped_at_creation: false,
         });
 
-        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 sort_buffer = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some(format!("request_sort_buffer_{N}").as_str()),
+            size: (size_of::<u32>() * index_buffer_count) as u64,
+            usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
+            mapped_at_creation: false,
+        });
 
         let bindgroup_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
             label: Some("request_sort_pipeline_bing_group_layout"),
@@ -160,7 +153,7 @@ where
                 },
                 BindGroupEntry {
                     binding: 2,
-                    resource: request_count_buffer.as_entire_binding(),
+                    resource: index_buffer.as_entire_binding(),
                 },
                 BindGroupEntry {
                     binding: 3,
@@ -168,7 +161,7 @@ where
                 },
                 BindGroupEntry {
                     binding: 4,
-                    resource: element_sort_count_buffer.as_entire_binding(),
+                    resource: request_count.as_entire_binding(),
                 },
             ],
         });
@@ -188,24 +181,13 @@ where
                     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; 
+                            @group(0) @binding(0) var<storage, read_write> request_buffer: array<u32>;
+                            @group(0) @binding(1) var<storage, read_write> sort_buffer: array<u32>; 
+                            @group(0) @binding(2) var<storage, read_write> index_buffer: array<u32>; 
+                            @group(0) @binding(3) var<storage, read_write> indirect_count: vec3<u32>; 
+                            @group(0) @binding(4) var<storage, read_write> request_count: u32; 
 
                             fn big_fusion(index: u32, phase_size: u32) -> vec2<u32>
                             {{
@@ -234,14 +216,14 @@ where
 
 
                             @compute
-                            @workgroup_size(64)
+                            @workgroup_size(64, 1, 1)
                             fn main(
                                     @builtin(global_invocation_id) global_invocation_id: vec3<u32>
                                 )
                             {{
 
                                  let index = global_invocation_id.x;
-                                 let total = arrayLength(&sort_indirection);
+                                 //let total = arrayLength(&sort_indirection);
 
                                  let sub_phase = (both_phase >> 16) & 0xFFFF;
                                  let phase = both_phase & 0xFFFF;
@@ -252,18 +234,30 @@ where
                                  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))
+                                      if(next_phase_width >= (request_count * 2) && phase_total_width >= (request_count * 2))
                                       {{
-                                        indirect_count = vec3<u32>(0);
+                                        //indirect_count = vec3<u32>(0);
                                       }}
                                  }}
 
                                  if(phase == 0)
                                  {{
+
+                                    // Phase 0. Index buffer contains indicies of compacted buffer.
+                                    // We can perform compaction on the data
                                     let a = index * 2;
                                     let b = a + 1;
-                                    sort_indirection[a].child = a % {children_count};
-                                    sort_indirection[b].child = b % {children_count};
+
+                                    if(a < request_count) 
+                                    {{
+                                         let compact_index_a = index_buffer[a];
+                                         sort_buffer[a] = request_buffer[compact_index_a];
+                                    }}
+                                    if(b < request_count) 
+                                    {{
+                                         let compact_index_b = index_buffer[b];
+                                         sort_buffer[b] = request_buffer[compact_index_b];
+                                    }}
                                  }}
 
                                  var swap_indices = vec2<u32>(0, 0);
@@ -280,180 +274,24 @@ where
                                  // Do swap
 
                                  //if(swap_indices.y >= request_count_round_up)
-                                 if(swap_indices.y >= element_sort_count)
+                                 if(swap_indices.y >= request_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];
+                                 let av = sort_buffer[swap_indices.x];
+                                 let bv = sort_buffer[swap_indices.y];
 
                                  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 sort_temp = sort_buffer[swap_indices.x];
+                                    sort_buffer[swap_indices.x] = sort_buffer[swap_indices.y];
+                                    sort_buffer[swap_indices.y] = sort_temp;
 
-        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;
-                                    }}
+                                    let index_temp = index_buffer[swap_indices.x];
+                                    index_buffer[swap_indices.x] = index_buffer[swap_indices.y];
+                                    index_buffer[swap_indices.y] = index_temp;
                                  }}
                             }}
                         ")
@@ -478,13 +316,14 @@ where
                                 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; 
+                            @group(0) @binding(0) var<storage, read_write> request_buffer: array<u32>;
+                            @group(0) @binding(1) var<storage, read_write> sort_buffer: array<u32>; 
+                            @group(0) @binding(2) var<storage, read_write> index_buffer: array<u32>; 
+                            @group(0) @binding(3) var<storage, read_write> indirect_count: vec3<u32>; 
+                            @group(0) @binding(4) var<storage, read_write> request_count: u32; 
 
                             @compute
-                            @workgroup_size(64)
+                            @workgroup_size(256, 1, 1)
                             fn main(
                                     @builtin(global_invocation_id) global_invocation_id: vec3<u32>
                                 )
@@ -494,20 +333,44 @@ where
                                     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;
-                                    }}
+                                    request_buffer[index] = 0;
+                            }}
+                        ")
+                        .into(),
+                    ),
+                }),
+                entry_point: Some("main"),
+                compilation_options: Default::default(),
+                cache: None,
+            });
+
+        let compute_indirect_pipeline =
+            device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+                label: Some("compute_indirect_pipeline"),
+                layout: Some(pipeline_layouts),
+                module: &device.create_shader_module(wgpu::ShaderModuleDescriptor {
+                    label: Some("clear_chunk_requests shader module"),
+                    source: wgpu::ShaderSource::Wgsl(
+                        format!("
+                            @group(0) @binding(0) var<storage, read_write> request_buffer: array<u32>;
+                            @group(0) @binding(1) var<storage, read_write> sort_buffer: array<u32>; 
+                            @group(0) @binding(2) var<storage, read_write> index_buffer: array<u32>; 
+                            @group(0) @binding(3) var<storage, read_write> indirect_count: vec3<u32>; 
+                            @group(0) @binding(4) var<storage, read_write> request_count: u32; 
+
+                            @compute
+                            @workgroup_size(1, 1, 1)
+                            fn main(
+                                    @builtin(global_invocation_id) global_invocation_id: vec3<u32>
+                                )
+                            {{
+                                 let invocation_count = (request_count / 2)
+                                     + select(u32(0), u32(1), request_count % 2 != 0);
+
+                                 let workgroup_count = (invocation_count / 64)
+                                     + select(u32(0), u32(1), invocation_count % 64 != 0);
+
+                                 indirect_count = vec3<u32>(workgroup_count, 1, 1);
                             }}
                         ")
                         .into(),
@@ -519,17 +382,17 @@ where
             });
 
         Self {
+            buffer_compactor: NonZeroBufferCompactor::new(&device, index_buffer_count),
             cache_size,
             request_buffer,
-            request_count_buffer,
-            indirect_count,
-            indirect_count_storage,
-            element_sort_count_buffer,
-            running_sum_pipeline,
+            index_buffer,
             sort_buffer,
+            indirect_count,
+            request_count,
+            indirect_count_storage,
+            compute_indirect_pipeline,
             sort_pipeline,
             reset_pipeline,
-            compaction_pipeline,
             bindgroup,
             device,
         }
@@ -542,12 +405,12 @@ where
 
     pub fn request_count_buffer(&self) -> &Buffer
     {
-        &self.request_count_buffer
+        &self.request_count
     }
 
-    pub fn sort_buffer(&self) -> &Buffer
+    pub fn index_buffer(&self) -> &Buffer
     {
-        &self.sort_buffer
+        &self.index_buffer
     }
 
     pub fn reset_requests(&self, encoder: &mut CommandEncoder)
@@ -560,57 +423,46 @@ where
         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);
+        let shader_invocations = self.cache_size * N * N * N; // one invocation per element
+        let workgroup_invocations = shader_invocations.div_ceil(256);
         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,
-            });
+        // Compact list
+        self.buffer_compactor.compact_buffer(
+            &self.device,
+            encoder,
+            &self.request_buffer,
+            &self.index_buffer,
+        );
 
-        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);
+        self.buffer_compactor
+            .copy_count_to_buffer(encoder, &self.request_count);
 
-        // = Perform list compaction
+        let mut compute_indirect_pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
+            label: Some("compute_indirect_pass"),
+            timestamp_writes: None,
+        });
 
-        // == Running sum
-        compaction_compute_pass.set_pipeline(&self.running_sum_pipeline);
+        compute_indirect_pass.set_bind_group(0, Some(&self.bindgroup), &[]);
 
-        // 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);
-        }
+        compute_indirect_pass.set_pipeline(&self.compute_indirect_pipeline);
+        compute_indirect_pass.dispatch_workgroups(1, 1, 1);
+        drop(compute_indirect_pass);
 
-        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);
-        }
+        // Now the indirect count storate buffer contains the workgroup count for the sort
+        // Copy that to the indirect buffer
+        encoder.copy_buffer_to_buffer(
+            &self.indirect_count_storage,
+            0,
+            &self.indirect_count,
+            0,
+            Some(size_of::<wgpu::util::DispatchIndirectArgs>() as u64),
+        );
 
-        // == 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
+        // Sort
         let sort_element_count = self.cache_size * N * N * N;
         let phases_upper_bound = sort_element_count.next_power_of_two().ilog2();