octagonal/vxls
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

feature: add a producer_interface struct to cleanly represent the $ interface

Browse Source
This commit is contained in:
octagonal
2026-09-05 21:06:38 +02:00
parent a241b7fd83
commit f8087c3a2c
3 changed files with 519 additions and 322 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/main.rs
+80 -110
View File
@@ -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<Mutex<VoxelCache<4>>>,
cache_interface: Arc<CacheProducerInterface<4>>,
terrain_generator: Arc<TerrainGenerator<4>>,
chunk_pos_map: Arc<HashMap<u32, (usize, usize, usize)>>,
instance_buffer: Buffer,
@@ -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(&timestamp_query, 0..2, &timestamp_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<CacheNodeRequest> =
let cache_node_requests: Vec<CacheRequest> =
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::<Vec<_>>()
.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()]);
}
},
src/voxel_cache.rs
+70 -212
View File
@@ -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<const N: usize>
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::<u32>() as u64,
usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
mapped_at_creation: false,
});
let request_interface_bind_group_layout = CacheProducerInterface::<N>::request_side_bind_group_layout(&device);
// Rendering bing group layouts
let voxel_cache_render_bind_group_layout =
@@ -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<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(0) @binding(7) var<storage, read_write> total_request_count: u32;
@group(1) @binding(0) var<storage, read_write> structure_table_pointers: array<u32>;
@group(1) @binding(1) var<storage, read_write> structure_table_sorted_requests: array<SortedRequestsElement>;
@group(1) @binding(2) var<storage, read_write> structure_table_request_count: u32;
");
let write_requests_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<storage, read_write> requests: array<CacheNodeRequest>;
struct CacheInterfaceRequestWb
{{
node_index: u32,
child_index: u32
}}
@group(2) @binding(0) var<storage, read_write> requests: array<CacheInterfaceRequest>;
@group(2) @binding(1) var<storage, read_write> requests_wb: array<CacheInterfaceRequestWb>;
@group(2) @binding(2) var<storage, read_write> request_count: u32;
@group(2) @binding(3) var<storage, read> structure_nodes: array<StructurePoolElement>;
@group(2) @binding(4) var<storage, read> color_nodes: array<ColorPoolElement>;
@group(2) @binding(5) var<storage, read> locations: array<LocationPoolElement>;
@compute
@workgroup_size(64)
@@ -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::<CachingPipelineImmediates>() as u32, // Current frame timestamp
});
@@ -611,10 +524,26 @@ where
var<immediate> parameters: CachingPipelineImmediates;
@group(2) @binding(0) var<storage, read> structure_nodes: array<StructurePoolElement>;
@group(2) @binding(1) var<storage, read> color_nodes: array<ColorPoolElement>;
@group(2) @binding(2) var<storage, read> locations: array<LocationPoolElement>;
@group(2) @binding(3) var<storage, read> destinations: array<DestinationElement>;
struct CacheInterfaceRequest
{{
structure_id: u32,
locator: u32,
child_index: u32
}}
struct CacheInterfaceRequestWb
{{
node_index: u32,
child_index: u32
}}
@group(2) @binding(0) var<storage, read_write> requests: array<CacheInterfaceRequest>;
@group(2) @binding(1) var<storage, read_write> requests_wb: array<CacheInterfaceRequestWb>;
@group(2) @binding(2) var<storage, read_write> request_count: u32;
@group(2) @binding(3) var<storage, read_write> structure_nodes: array<StructurePoolElement>;
@group(2) @binding(4) var<storage, read_write> color_nodes: array<ColorPoolElement>;
@group(2) @binding(5) var<storage, read_write> locations: array<LocationPoolElement>;
@compute
@workgroup_size(64)
@@ -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<N>)
{
// 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::<CacheNodeRequest>() 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<N>)
{
// Each of the entry of the cache insertion, is matched with the usage list to evict
let mut cache_insertion_pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
@@ -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::<StructurePoolElement<N>>();
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::<StructurePoolElement<N>>();
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.");
}
}
src/voxel_cache/producer_interface.rs
+369
View File
@@ -0,0 +1,369 @@
// The cache emits requests, and receives answer to these requests
// The cache producer interface provides the necessary buffers to contains
// these, so that use applications can fullfill the requests
use bytemuck::Pod;
use bytemuck::Zeroable;
use wgpu::BindGroup;
use wgpu::BindGroupEntry;
use wgpu::BindGroupLayout;
use wgpu::BindGroupLayoutEntry;
use wgpu::Buffer;
use wgpu::BufferUsages;
use wgpu::CommandEncoder;
use wgpu::Device;
use wgpu::Queue;
use wgpu::util::DownloadBuffer;
use crate::voxel_cache::data::ColorPoolElement;
use crate::voxel_cache::data::LocationPoolElement;
use crate::voxel_cache::data::StructurePoolElement;
#[derive(Clone, Copy, Zeroable, Pod)]
#[repr(C)]
pub struct CacheRequest
{
pub structure_id: u32,
pub locator: u32,
pub child_index: u32,
}
pub(super) struct CacheRequestWbInfo
{
node_index: u32,
child_index: u32,
}
pub struct CacheProducerInterface<const N: usize>
where
[(); N * N * N]:,
{
// Request emissions
pub(super) size: usize,
// Contains the specific node being requested
pub(super) requests: Buffer,
// Contains how many request have effectively been emitted
pub(super) request_count: Buffer,
// Contains on which node the request was done
pub(super) requests_wb_info: Buffer,
// Request fullfillment
pub(super) structure_nodes: Buffer,
pub(super) color_nodes: Buffer,
pub(super) location_nodes: Buffer,
pub(super) request_side_bind_group_layout: BindGroupLayout,
pub(super) request_side_bind_group: BindGroup,
pub(super) producer_side_bind_group_layout: BindGroupLayout,
pub(super) producer_side_bind_group: BindGroup,
}
impl<const N: usize> CacheProducerInterface<N>
where
[(); N * N * N]:,
{
pub fn new(size: usize, device: &Device) -> Self
{
let requests = device.create_buffer(&wgpu::BufferDescriptor {
label: Some("cache_interface_request_buffer"),
size: (size_of::<CacheRequest>() * size) as u64,
usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
mapped_at_creation: false,
});
let requests_wb_info = device.create_buffer(&wgpu::BufferDescriptor {
label: Some("cache_interface_writeback_info"),
size: (size_of::<CacheRequestWbInfo>() * size) as u64,
usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
mapped_at_creation: false,
});
let structure_nodes = device.create_buffer(&wgpu::BufferDescriptor {
label: Some("cache_interface_structure_nodes"),
size: (size_of::<StructurePoolElement<N>>() * size) as u64,
usage: BufferUsages::STORAGE | BufferUsages::COPY_DST,
mapped_at_creation: false,
});
let color_nodes = device.create_buffer(&wgpu::BufferDescriptor {
label: Some("cache_interface_structure_nodes"),
size: (size_of::<ColorPoolElement<N>>() * size) as u64,
usage: BufferUsages::STORAGE | BufferUsages::COPY_DST,
mapped_at_creation: false,
});
let location_nodes = device.create_buffer(&wgpu::BufferDescriptor {
label: Some("cache_interface_structure_nodes"),
size: (size_of::<LocationPoolElement>() * size) as u64,
usage: BufferUsages::STORAGE | BufferUsages::COPY_DST,
mapped_at_creation: false,
});
let request_count = device.create_buffer(&wgpu::BufferDescriptor {
label: Some("cache_interface_request_count"),
size: (size_of::<u32>()) as u64,
usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
mapped_at_creation: false,
});
let request_side_bind_group_layout = Self::request_side_bind_group_layout(device);
let producer_side_bind_group_layout =
device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
label: Some("cache_interface_producer_side_bind_group_layout"),
entries: &[
// Reqests
BindGroupLayoutEntry {
binding: 0,
visibility: wgpu::ShaderStages::COMPUTE,
ty: wgpu::BindingType::Buffer {
ty: wgpu::BufferBindingType::Storage { read_only: true },
has_dynamic_offset: false,
min_binding_size: None,
},
count: None,
},
// Request count
BindGroupLayoutEntry {
binding: 1,
visibility: wgpu::ShaderStages::COMPUTE,
ty: wgpu::BindingType::Buffer {
ty: wgpu::BufferBindingType::Storage { read_only: true },
has_dynamic_offset: false,
min_binding_size: None,
},
count: None,
},
// Structure nodes
BindGroupLayoutEntry {
binding: 2,
visibility: wgpu::ShaderStages::COMPUTE,
ty: wgpu::BindingType::Buffer {
ty: wgpu::BufferBindingType::Storage { read_only: false },
has_dynamic_offset: false,
min_binding_size: None,
},
count: None,
},
// Color nodes
BindGroupLayoutEntry {
binding: 3,
visibility: wgpu::ShaderStages::COMPUTE,
ty: wgpu::BindingType::Buffer {
ty: wgpu::BufferBindingType::Storage { read_only: false },
has_dynamic_offset: false,
min_binding_size: None,
},
count: None,
},
// Location nodes
BindGroupLayoutEntry {
binding: 4,
visibility: wgpu::ShaderStages::COMPUTE,
ty: wgpu::BindingType::Buffer {
ty: wgpu::BufferBindingType::Storage { read_only: false },
has_dynamic_offset: false,
min_binding_size: None,
},
count: None,
},
],
});
// Bind groups
let request_side_bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
label: Some("cache_interface_request_side_bind_group"),
layout: &request_side_bind_group_layout,
entries: &[
BindGroupEntry {
binding: 0,
resource: requests.as_entire_binding(),
},
BindGroupEntry {
binding: 1,
resource: requests_wb_info.as_entire_binding(),
},
BindGroupEntry {
binding: 2,
resource: request_count.as_entire_binding(),
},
BindGroupEntry {
binding: 3,
resource: structure_nodes.as_entire_binding(),
},
BindGroupEntry {
binding: 4,
resource: color_nodes.as_entire_binding(),
},
BindGroupEntry {
binding: 5,
resource: location_nodes.as_entire_binding(),
},
],
});
let producer_side_bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
label: Some("cache_interface_producer_side_bind_group"),
layout: &producer_side_bind_group_layout,
entries: &[
BindGroupEntry {
binding: 0,
resource: requests.as_entire_binding(),
},
BindGroupEntry {
binding: 1,
resource: request_count.as_entire_binding(),
},
BindGroupEntry {
binding: 2,
resource: structure_nodes.as_entire_binding(),
},
BindGroupEntry {
binding: 3,
resource: color_nodes.as_entire_binding(),
},
BindGroupEntry {
binding: 4,
resource: location_nodes.as_entire_binding(),
},
],
});
Self {
size,
requests,
request_count,
requests_wb_info,
structure_nodes,
color_nodes,
location_nodes,
request_side_bind_group_layout,
request_side_bind_group,
producer_side_bind_group_layout,
producer_side_bind_group,
}
}
pub fn requests_buffer(&self) -> &Buffer
{
&self.requests
}
pub fn structure_nodes_buffer(&self) -> &Buffer
{
&self.structure_nodes
}
pub fn color_nodes_buffer(&self) -> &Buffer
{
&self.color_nodes
}
pub fn location_nodes_buffer(&self) -> &Buffer
{
&self.location_nodes
}
pub fn total_request_count(&self, device: &Device, queue: &Queue) -> u32
{
let (tx, rx) = std::sync::mpsc::sync_channel(1);
wgpu::util::DownloadBuffer::read_buffer(
device,
queue,
&self.request_count.slice(0..),
move |download_buffer| {
let vec = download_buffer.unwrap().to_vec();
let value = u32::from_ne_bytes(std::array::from_fn(|i| vec[i]));
let _ = tx.try_send(value);
},
);
loop
{
if let Ok(value) = rx.try_recv()
{
return value;
}
device.poll(wgpu::wgt::PollType::Poll).unwrap();
}
//panic!("Could not retrieve total request count.");
}
pub fn request_side_bind_group_layout(device: &Device) -> BindGroupLayout
{
device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
label: Some("cache_interface_request_side_bind_group_layout"),
entries: &[
// Reqests
BindGroupLayoutEntry {
binding: 0,
visibility: wgpu::ShaderStages::COMPUTE,
ty: wgpu::BindingType::Buffer {
ty: wgpu::BufferBindingType::Storage { read_only: false },
has_dynamic_offset: false,
min_binding_size: None,
},
count: None,
},
// Requests writeback
BindGroupLayoutEntry {
binding: 1,
visibility: wgpu::ShaderStages::COMPUTE,
ty: wgpu::BindingType::Buffer {
ty: wgpu::BufferBindingType::Storage { read_only: false },
has_dynamic_offset: false,
min_binding_size: None,
},
count: None,
},
// Request count
BindGroupLayoutEntry {
binding: 2,
visibility: wgpu::ShaderStages::COMPUTE,
ty: wgpu::BindingType::Buffer {
ty: wgpu::BufferBindingType::Storage { read_only: false },
has_dynamic_offset: false,
min_binding_size: None,
},
count: None,
},
// Structure nodes
BindGroupLayoutEntry {
binding: 3,
visibility: wgpu::ShaderStages::COMPUTE,
ty: wgpu::BindingType::Buffer {
ty: wgpu::BufferBindingType::Storage { read_only: false },
has_dynamic_offset: false,
min_binding_size: None,
},
count: None,
},
// Color nodes
BindGroupLayoutEntry {
binding: 4,
visibility: wgpu::ShaderStages::COMPUTE,
ty: wgpu::BindingType::Buffer {
ty: wgpu::BufferBindingType::Storage { read_only: false },
has_dynamic_offset: false,
min_binding_size: None,
},
count: None,
},
// Location nodes
BindGroupLayoutEntry {
binding: 5,
visibility: wgpu::ShaderStages::COMPUTE,
ty: wgpu::BindingType::Buffer {
ty: wgpu::BufferBindingType::Storage { read_only: false },
has_dynamic_offset: false,
min_binding_size: None,
},
count: None,
},
],
})
}
}