diff --git a/src/main.rs b/src/main.rs index 5598674..1eb9efe 100644 --- a/src/main.rs +++ b/src/main.rs @@ -59,6 +59,8 @@ 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::producer_interface::CacheProducerInterface; +use crate::voxel_cache::producer_interface::CacheRequest; mod camera; mod egui_renderer; @@ -81,6 +83,7 @@ struct State pipeline: RenderPipeline, voxel_cache: Arc>>, + cache_interface: Arc>, terrain_generator: Arc>, chunk_pos_map: Arc>, instance_buffer: Buffer, @@ -156,20 +159,21 @@ impl State let egui_renderer = EguiRenderer::new(&device, surface_format, &window); let mut voxel_cache = VoxelCache::<4>::new(200_000, device.clone(), queue.clone()); + let cache_interface = CacheProducerInterface::new(1024, &device); - //let terrain_generator = TerrainGenerator::<4>::new(5, "vxls_height.tif", 0.2, "img.jpg"); + 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 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) @@ -280,6 +284,7 @@ impl State 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, @@ -380,6 +385,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, @@ -407,6 +413,16 @@ impl State } }; + let texture_view = surface_texture + .texture + .create_view(&wgpu::TextureViewDescriptor { + // Without add_srgb_suffix() the image we will be working with + // might not be "gamma correct". + format: Some(self.surface_format.add_srgb_suffix()), + ..Default::default() + }); + + // ~~ Ray-marching timestamp query setup ~~ let timestamp_query = self.device.create_query_set(&wgpu::QuerySetDescriptor { label: Some("timestamp_query_set"), ty: wgpu::QueryType::Timestamp, @@ -420,18 +436,8 @@ impl State mapped_at_creation: false, }); - let texture_view = surface_texture - .texture - .create_view(&wgpu::TextureViewDescriptor { - // Without add_srgb_suffix() the image we will be working with - // might not be "gamma correct". - format: Some(self.surface_format.add_srgb_suffix()), - ..Default::default() - }); - - // Renders a GREEN screen + // ~~ Main render pass ~~ let mut encoder = self.device.create_command_encoder(&Default::default()); - { let mut renderpass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { label: None, @@ -440,7 +446,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, }, })], @@ -478,24 +489,7 @@ impl State encoder.resolve_query_set(×tamp_query, 0..2, ×tamp_buffer, 0); } - let requests = self.device.create_buffer(&wgpu::BufferDescriptor { - label: Some("dummy_dumb_dinky_aaaahhh_buffer"), - size: 16 * 2000, - 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( @@ -547,7 +541,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( @@ -567,27 +561,7 @@ 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); - if !self - .camera - .pressed_keyset - .contains(&winit::keyboard::KeyCode::KeyF) - && self.insertion_debounce - { - self.insertion_debounce = false; - } - - // if (self - // .camera - // .pressed_keyset - // .contains(&winit::keyboard::KeyCode::KeyF)) - // && !self.insertion_debounce - // { - - // Report frame time + // ~~ Get Ray-marching timestamps, report time ~~ let cloned_rm_time = self.rm_time.clone(); let cloned_queue = self.queue.clone(); DownloadBuffer::read_buffer( @@ -602,30 +576,52 @@ impl State }, ); + // ~~ Do cache managment + if !self + .camera + .pressed_keyset + .contains(&winit::keyboard::KeyCode::KeyF) + { + self.voxel_cache + .lock() + .cache_post_render(&mut encoder, &self.cache_interface); + } + + // ~~ Submit command buffer ~~ + 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 = 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 = + let cache_node_requests: Vec = bytemuck::pod_collect_to_vec(&buffer.unwrap()); let generator = cloned_generator; @@ -634,7 +630,6 @@ impl State let mut structure_nodes = vec![]; let mut color_nodes: Vec<[ColorBytes; 64]> = vec![]; let mut location_nodes = vec![]; - let mut destinations = vec![]; cache_node_requests .par_iter() @@ -646,7 +641,7 @@ impl State let location; let node; - if request.child_index == u32::MAX + if request.locator == 0 { // Produce root node node = @@ -660,9 +655,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(); @@ -692,15 +686,7 @@ impl State }; } - ( - node.structure, - node.colors, - location, - DestinationElement { - node: request.node_index, - child: request.child_index, - }, - ) + (node.structure, node.colors, location) }) .collect::>() .into_iter() @@ -708,46 +694,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()]); } }, diff --git a/src/voxel_cache.rs b/src/voxel_cache.rs index d4ef929..7b1d1c8 100644 --- a/src/voxel_cache.rs +++ b/src/voxel_cache.rs @@ -1,12 +1,13 @@ use bytemuck::Zeroable; -use wgpu::{BindGroup, BindGroupLayout, Buffer, BufferUsages, ComputePipeline, Device, Queue, ShaderStages, CommandEncoder}; +use wgpu::{BindGroup, BindGroupLayout, Buffer, BufferUsages, CommandEncoder, ComputePipeline, Device, Queue, ShaderStages}; -use crate::voxel_cache::{data::{CacheNodeRequest, CacheResponse, ColorPoolElement, LocationPoolElement, StructurePoolElement}, request_buffer::RequestBuffer, structure_table::StructureTable, usage_buffer::UsageBuffer}; +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 data; @@ -28,12 +29,9 @@ pub struct VoxelCache voxel_cache_render_bind_group_layout: BindGroupLayout, // Writing requests - request_write_count_buffer: Buffer, - request_write_bindgroup_layout_user: BindGroupLayout, request_write_pipeline: ComputePipeline, // User response -> caching - caching_user_bindgroup_layout: BindGroupLayout, caching_pipeline: ComputePipeline, // Invalidation @@ -87,12 +85,7 @@ where let request_buffer = RequestBuffer::new(cache_size, device.clone()); let usage_buffer = UsageBuffer::new(cache_size, device.clone()); - let request_write_count_buffer = device.create_buffer(&wgpu::BufferDescriptor { - label: Some("total_request_count_buffer"), - size: size_of::() as u64, - usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC, - mapped_at_creation: false, - }); + let request_interface_bind_group_layout = CacheProducerInterface::::request_side_bind_group_layout(&device); // Rendering bing group layouts let voxel_cache_render_bind_group_layout = @@ -263,17 +256,6 @@ where }, count: None, }, - // total request count - wgpu::BindGroupLayoutEntry { - binding: 7, - visibility: ShaderStages::COMPUTE, - ty: wgpu::BindingType::Buffer { - ty: wgpu::BufferBindingType::Storage { read_only: false }, - has_dynamic_offset: false, - min_binding_size: None, - }, - count: None, - }, ], }); @@ -316,11 +298,6 @@ where binding: 6, resource: request_buffer.request_count_buffer().as_entire_binding(), }, - // Total request count - wgpu::BindGroupEntry { - binding: 7, - resource: request_write_count_buffer.as_entire_binding(), - }, ], }); @@ -396,39 +373,19 @@ where @group(0) @binding(4) var lru_list: array; @group(0) @binding(5) var usage_buffer: array; @group(0) @binding(6) var pools_request_count: u32; - @group(0) @binding(7) var total_request_count: u32; @group(1) @binding(0) var structure_table_pointers: array; @group(1) @binding(1) var structure_table_sorted_requests: array; @group(1) @binding(2) var structure_table_request_count: u32; "); - - let write_requests_bindgroup_layout_user = - device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { - label: Some("write_requests_bindgroup_layout_user"), - entries: &[ - // User target - 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 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(&write_requests_bindgroup_layout_user), + Some(&request_interface_bind_group_layout), ], immediate_size: 0, }); @@ -444,18 +401,26 @@ where format!(" {wgsl_bindings} - struct CacheNodeRequest + struct CacheInterfaceRequest {{ - // Requested ressource structure_id: u32, - structure_locator: u32, - - // Write back info - node_index: u32, - child_index: u32, + locator: u32, + child_index: u32 }} - @group(2) @binding(0) var requests: array; + struct CacheInterfaceRequestWb + {{ + node_index: u32, + child_index: u32 + }} + + @group(2) @binding(0) var requests: array; + @group(2) @binding(1) var requests_wb: array; + @group(2) @binding(2) var request_count: u32; + + @group(2) @binding(3) var structure_nodes: array; + @group(2) @binding(4) var color_nodes: array; + @group(2) @binding(5) var locations: array; @compute @workgroup_size(64) @@ -467,38 +432,45 @@ where // and cache domain var index = global_invocation_id.x; let max_requests_count = min(arrayLength(&requests), arrayLength(&lru_list)); - total_request_count = min(max_requests_count, pools_request_count + structure_table_request_count); - if(index >= total_request_count) + 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: CacheNodeRequest; + var request: CacheInterfaceRequest; // Index in structure table IS structure id request.structure_id = structure_table_sorted_requests[index].node; - request.structure_locator = 0; // Root request -> locator 0 - request.node_index = request.structure_id; + 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; - var request: CacheNodeRequest; // Request location comes from location pool + var request: CacheInterfaceRequest; request.structure_id = location_pool[sorted_requests[pool_index].node].structure_id; - request.structure_locator = location_pool[sorted_requests[pool_index].node].structure_locator; - - request.node_index = sorted_requests[pool_index].node; + 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; }} " @@ -515,72 +487,13 @@ where // Pipeline that takes user fulling ~some~ requests // and writes them to the cache based on the eviction list - // Caching pipeline binds - // - Cache data - // - // - User response : - // - response structure nodes - // - response color nodes - // - response locations - // - response destinations - let caching_user_bindgroup_layout = - device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { - label: Some("caching_user_bindgroup_layout"), - entries: &[ - // structure nodes - wgpu::BindGroupLayoutEntry { - binding: 0, - visibility: ShaderStages::COMPUTE, - ty: wgpu::BindingType::Buffer { - ty: wgpu::BufferBindingType::Storage { read_only: true }, - has_dynamic_offset: false, - min_binding_size: None, - }, - count: None, - }, - // response color nodes - wgpu::BindGroupLayoutEntry { - binding: 1, - visibility: ShaderStages::COMPUTE, - ty: wgpu::BindingType::Buffer { - ty: wgpu::BufferBindingType::Storage { read_only: true }, - has_dynamic_offset: false, - min_binding_size: None, - }, - count: None, - }, - // response locations - wgpu::BindGroupLayoutEntry { - binding: 2, - visibility: ShaderStages::COMPUTE, - ty: wgpu::BindingType::Buffer { - ty: wgpu::BufferBindingType::Storage { read_only: true }, - has_dynamic_offset: false, - min_binding_size: None, - }, - count: None, - }, - // response destinations - 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, - }, - ], - }); - 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(&caching_user_bindgroup_layout), + Some(&request_interface_bind_group_layout), ], immediate_size: size_of::() as u32, // Current frame timestamp }); @@ -611,10 +524,26 @@ where var parameters: CachingPipelineImmediates; - @group(2) @binding(0) var structure_nodes: array; - @group(2) @binding(1) var color_nodes: array; - @group(2) @binding(2) var locations: array; - @group(2) @binding(3) var destinations: array; + struct CacheInterfaceRequest + {{ + structure_id: u32, + locator: u32, + child_index: u32 + }} + + struct CacheInterfaceRequestWb + {{ + node_index: u32, + child_index: u32 + }} + + @group(2) @binding(0) var requests: array; + @group(2) @binding(1) var requests_wb: array; + @group(2) @binding(2) var request_count: u32; + + @group(2) @binding(3) var structure_nodes: array; + @group(2) @binding(4) var color_nodes: array; + @group(2) @binding(5) var locations: array; @compute @workgroup_size(64) @@ -653,12 +582,12 @@ where // Point parent to new page let new_pointer = (1 << 31) | (1 << 30) | overwritten_element; - if(destinations[index].child == 0xFFFFFFFF) + if(requests_wb[index].child_index == 0xFFFFFFFF) {{ - structure_table_pointers[destinations[index].node] = new_pointer; - }}else if usage_buffer[destinations[index].node] != parameters.frame_timestamp + 1 + 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[destinations[index].node].pointers[destinations[index].child] = new_pointer; + structure_pool[requests_wb[index].node_index].pointers[requests_wb[index].child_index] = new_pointer; }} }} @@ -769,12 +698,9 @@ where voxel_cache_render_bind_group_layout, // Write requests stage - request_write_bindgroup_layout_user: write_requests_bindgroup_layout_user, - request_write_count_buffer, request_write_pipeline: write_requests, // Caching operation - caching_user_bindgroup_layout, caching_pipeline, // Invalidation @@ -860,7 +786,7 @@ where }) } - pub fn cache_post_render(&mut self, encoder: &mut CommandEncoder, request_target: Buffer) + pub fn cache_post_render(&mut self, encoder: &mut CommandEncoder, request_interface: &CacheProducerInterface) { // Sort requests, reset request buffers to count requests self.usage_buffer.sort_usage(encoder); @@ -871,18 +797,6 @@ where .sort_requests(encoder); // Sorted requests are now in the group - let user_target_bindgroup = self.device.create_bind_group(&wgpu::BindGroupDescriptor { - label: Some("write_requests_user_bindgroup"), - layout: &self.request_write_bindgroup_layout_user, - entries: &[ - // user target - wgpu::BindGroupEntry { - binding: 0, - resource: request_target.as_entire_binding(), - }, - ], - }); - let mut write_requests_pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { label: Some("write_requests_pass"), timestamp_writes: None, @@ -890,11 +804,11 @@ where 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(&user_target_bindgroup), &[]); + 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_target.size() / size_of::() as u64; + 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); @@ -910,7 +824,7 @@ where self.usage_buffer.timestamp() } - pub fn cache_insert(&mut self, encoder: &mut CommandEncoder, insertion: &CacheResponse) + pub fn cache_insert(&mut self, encoder: &mut CommandEncoder, cache_interface: &CacheProducerInterface) { // 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 { @@ -918,35 +832,6 @@ where timestamp_writes: None, }); - let user_caching_bindgroup = self.device.create_bind_group(&wgpu::BindGroupDescriptor - { - label: Some("user_caching_bind_group"), - layout: &self.caching_user_bindgroup_layout, - entries: & - [ - wgpu::BindGroupEntry - { - binding: 0, - resource: insertion.structure_nodes.as_entire_binding(), - }, - wgpu::BindGroupEntry - { - binding: 1, - resource: insertion.color_nodes.as_entire_binding(), - }, - wgpu::BindGroupEntry - { - binding: 2, - resource: insertion.locations.as_entire_binding(), - }, - wgpu::BindGroupEntry - { - binding: 3, - resource: insertion.parents.as_entire_binding(), - } - ], - }); - // Copy insertions to evicted lru // Point parents to children let structure_table_bind_group = self.structure_table_bind_group(); @@ -956,7 +841,7 @@ where 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(&user_caching_bindgroup), &[]); + cache_insertion_pass.set_bind_group(2, Some(&cache_interface.request_side_bind_group), &[]); cache_insertion_pass.set_immediates( 0, bytemuck::bytes_of(&CachingPipelineImmediates { @@ -966,8 +851,7 @@ where ); // Compute dispatch amounts - let shader_invocation_count = - insertion.structure_nodes.size() as usize / size_of::>(); + 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); @@ -990,7 +874,7 @@ where // ~~~ Write new pointers ~~~ { cache_insertion_pass.set_pipeline(&self.caching_pipeline); - cache_insertion_pass.set_bind_group(2, Some(&user_caching_bindgroup), &[]); + cache_insertion_pass.set_bind_group(2, Some(&cache_interface.request_side_bind_group), &[]); cache_insertion_pass.set_immediates( 0, bytemuck::bytes_of(&CachingPipelineImmediates { @@ -1000,8 +884,7 @@ where ); // Compute dispatch amounts - let shader_invocation_count = - insertion.structure_nodes.size() as usize / size_of::>(); + 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); @@ -1014,29 +897,4 @@ where self.usage_buffer.next_frame(); } - pub fn total_request_count(&self) -> u32 - { - let (tx, rx) = std::sync::mpsc::sync_channel(1); - wgpu::util::DownloadBuffer::read_buffer( - &self.device, - &self.queue, - &self.request_write_count_buffer.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; - } - self.device.poll(wgpu::wgt::PollType::Poll).unwrap(); - } - - //panic!("Could not retrieve total request count."); - } } diff --git a/src/voxel_cache/producer_interface.rs b/src/voxel_cache/producer_interface.rs new file mode 100644 index 0000000..cbe79a3 --- /dev/null +++ b/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 +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 CacheProducerInterface +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::() * 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::() * 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::>() * 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::>() * 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::() * 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::()) 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, + }, + ], + }) + } +}