Smart compaction goes brrrrr
This commit is contained in:
@@ -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;
|
||||
}
|
||||
|
||||
Binary file not shown.
+20
-97
@@ -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(×tamp_query, 0);
|
||||
{
|
||||
let mut renderpass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
|
||||
label: Some("compute_render_pass"),
|
||||
timestamp_writes: Some(wgpu::ComputePassTimestampWrites {
|
||||
query_set: ×tamp_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(×tamp_query, 0..2, ×tamp_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(×tamp_query, 1);
|
||||
encoder.resolve_query_set(×tamp_query, 0..4, ×tamp_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;
|
||||
},
|
||||
);
|
||||
|
||||
|
||||
+14
-11
@@ -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
|
||||
{
|
||||
|
||||
+143
-291
@@ -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();
|
||||
|
||||
|
||||
Reference in New Issue
Block a user