diff --git a/Cargo.toml b/Cargo.toml index a460b71..12fa46f 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -1,6 +1,6 @@ [package] name = "shocovox-rs" -version = "0.4.2" +version = "0.5.0" edition = "2021" authors = ["Dávid Tóth "] license = "MIT OR Apache-2.0" diff --git a/assets/shaders/viewport_render.wgsl b/assets/shaders/viewport_render.wgsl index 6bbc932..55d4f09 100644 --- a/assets/shaders/viewport_render.wgsl +++ b/assets/shaders/viewport_render.wgsl @@ -283,6 +283,10 @@ fn set_brick_used(brick_index: u32) { } } +fn request_node(node: u32, child_octant: u32){ + +} + //crate::spatial::math::step_octant fn step_octant(octant: u32, step: ptr) -> u32 { return ( @@ -404,7 +408,7 @@ fn probe_brick( let brick_index = node_children[((leaf_node_key * 8) + brick_octant)]; //set_brick_used(brick_index); // +++ DEBUG +++ - if(0 == metadata[arrayLength(&metadata) - 1]){ + if(0 == debug_interface){ set_brick_used(brick_index); } if(0 != ((0x01u << (8 + brick_octant)) & metadata[leaf_node_key])) { // brick is not empty @@ -804,6 +808,12 @@ var voxels: array; @group(1) @binding(5) var color_palette: array; +// +++ DEBUG +++ +@group(1) @binding(6) +var debug_interface: u32; +// --- DEBUG --- + + @compute @workgroup_size(8, 8, 1) fn update( @builtin(global_invocation_id) invocation_id: vec3, diff --git a/examples/dot_cube.rs b/examples/dot_cube.rs index b4d02ad..12ce509 100644 --- a/examples/dot_cube.rs +++ b/examples/dot_cube.rs @@ -9,10 +9,8 @@ use iyes_perf_ui::{ #[cfg(feature = "bevy_wgpu")] use shocovox_rs::octree::{ - raytracing::{ - bevy::create_viewing_glass, ShocoVoxRenderPlugin, ShocoVoxViewingGlass, Viewport, - }, - Albedo, V3c, + raytracing::{OctreeGPUView, Ray, ShocoVoxRenderPlugin, Viewport}, + Albedo, Octree, V3c, VoxelData, }; #[cfg(feature = "bevy_wgpu")] @@ -24,6 +22,15 @@ const BRICK_DIMENSION: usize = 2; #[cfg(feature = "bevy_wgpu")] const TREE_SIZE: u32 = 16; +#[cfg(feature = "bevy_wgpu")] +#[derive(Resource)] +struct TreeResource +where + T: Default + Clone + PartialEq + VoxelData, +{ + tree: Octree, +} + #[cfg(feature = "bevy_wgpu")] fn main() { App::new() @@ -109,14 +116,14 @@ fn setup(mut commands: Commands, images: ResMut>) { } } - let render_data = tree.create_bevy_view(); - let viewing_glass = create_viewing_glass( - &Viewport { + let output_texture = tree.create_new_gpu_view( + Viewport { origin, direction: (V3c::new(0., 0., 0.) - origin).normalized(), w_h_fov: V3c::new(10., 10., 3.), }, DISPLAY_RESOLUTION, + &mut commands, images, ); @@ -130,13 +137,11 @@ fn setup(mut commands: Commands, images: ResMut>) { custom_size: Some(Vec2::new(1024., 768.)), ..default() }, - texture: viewing_glass.output_texture.clone(), + texture: output_texture, ..default() }); commands.spawn(Camera2dBundle::default()); - commands.insert_resource(render_data); - commands.insert_resource(viewing_glass); - + commands.insert_resource(TreeResource { tree }); commands.spawn(( PerfUiRoot::default(), PerfUiEntryFPS { @@ -163,31 +168,34 @@ struct DomePosition { yaw: f32, roll: f32, } - #[cfg(feature = "bevy_wgpu")] -fn rotate_camera( - angles_query: Query<&mut DomePosition>, - mut viewing_glass: ResMut, -) { +fn rotate_camera(angles_query: Query<&mut DomePosition>, mut tree_view: ResMut) { let (yaw, roll) = (angles_query.single().yaw, angles_query.single().roll); - let radius = angles_query.single().radius; - viewing_glass.viewport.origin = V3c::new( + + tree_view.viewing_glass.viewport.origin = V3c::new( radius / 2. + yaw.sin() * radius, radius + roll.sin() * radius * 2., radius / 2. + yaw.cos() * radius, ); - viewing_glass.viewport.direction = - (V3c::unit(radius / 2.) - viewing_glass.viewport.origin).normalized(); + tree_view.viewing_glass.viewport.direction = + (V3c::unit(radius / 2.) - tree_view.viewing_glass.viewport.origin).normalized(); } #[cfg(feature = "bevy_wgpu")] fn handle_zoom( keys: Res>, - mut viewing_glass: ResMut, + mut tree_view: ResMut, mut angles_query: Query<&mut DomePosition>, + tree: Res>, ) { - const ADDITION: f32 = 0.02; + if keys.pressed(KeyCode::Delete) { + tree_view.do_the_thing = true; + } + // if tree_view.is_changed() { + // println!("changed!! --> {:?}", tree_view.read_back); + // } + const ADDITION: f32 = 0.05; let angle_update_fn = |angle, delta| -> f32 { let new_angle = angle + delta; if new_angle < 360. { @@ -196,11 +204,75 @@ fn handle_zoom( 0. } }; + if keys.pressed(KeyCode::Tab) { + // Render the current view with CPU + let viewport_up_direction = V3c::new(0., 1., 0.); + let viewport_right_direction = viewport_up_direction + .cross(tree_view.viewing_glass.viewport.direction) + .normalized(); + let pixel_width = + tree_view.viewing_glass.viewport.w_h_fov.x as f32 / DISPLAY_RESOLUTION[0] as f32; + let pixel_height = + tree_view.viewing_glass.viewport.w_h_fov.y as f32 / DISPLAY_RESOLUTION[1] as f32; + let viewport_bottom_left = tree_view.viewing_glass.viewport.origin + + (tree_view.viewing_glass.viewport.direction + * tree_view.viewing_glass.viewport.w_h_fov.z) + - (viewport_up_direction * (tree_view.viewing_glass.viewport.w_h_fov.y / 2.)) + - (viewport_right_direction * (tree_view.viewing_glass.viewport.w_h_fov.x / 2.)); + + // define light + let diffuse_light_normal = V3c::new(0., -1., 1.).normalized(); + + use image::ImageBuffer; + use image::Rgb; + let mut img = ImageBuffer::new(DISPLAY_RESOLUTION[0], DISPLAY_RESOLUTION[1]); + + // cast each ray for a hit + for x in 0..DISPLAY_RESOLUTION[0] { + for y in 0..DISPLAY_RESOLUTION[1] { + let actual_y_in_image = DISPLAY_RESOLUTION[1] - y - 1; + //from the origin of the camera to the current point of the viewport + let glass_point = viewport_bottom_left + + viewport_right_direction * x as f32 * pixel_width + + viewport_up_direction * y as f32 * pixel_height; + let ray = Ray { + origin: tree_view.viewing_glass.viewport.origin, + direction: (glass_point - tree_view.viewing_glass.viewport.origin).normalized(), + }; + + use std::io::Write; + std::io::stdout().flush().ok().unwrap(); + + if let Some(hit) = tree.tree.get_by_ray(&ray) { + let (data, _, normal) = hit; + //Because both vector should be normalized, the dot product should be 1*1*cos(angle) + //That means it is in range -1, +1, which should be accounted for + let diffuse_light_strength = + 1. - (normal.dot(&diffuse_light_normal) / 2. + 0.5); + img.put_pixel( + x, + actual_y_in_image, + Rgb([ + (data.r as f32 * diffuse_light_strength) as u8, + (data.g as f32 * diffuse_light_strength) as u8, + (data.b as f32 * diffuse_light_strength) as u8, + ]), + ); + } else { + img.put_pixel(x, actual_y_in_image, Rgb([128, 128, 128])); + } + } + } + + img.save("example_junk_cpu_render.png").ok().unwrap(); + } + let multiplier = if keys.pressed(KeyCode::ShiftLeft) { 10.0 // Doesn't have any effect?! } else { 1.0 }; + if keys.pressed(KeyCode::ArrowUp) { angles_query.single_mut().roll = angle_update_fn(angles_query.single().roll, ADDITION); } @@ -222,12 +294,12 @@ fn handle_zoom( angles_query.single_mut().radius *= 1. + 0.02 * multiplier; } if keys.pressed(KeyCode::Home) { - viewing_glass.viewport.w_h_fov.x *= 1. + 0.09 * multiplier; - viewing_glass.viewport.w_h_fov.y *= 1. + 0.09 * multiplier; + tree_view.viewing_glass.viewport.w_h_fov.x *= 1. + 0.09 * multiplier; + tree_view.viewing_glass.viewport.w_h_fov.y *= 1. + 0.09 * multiplier; } if keys.pressed(KeyCode::End) { - viewing_glass.viewport.w_h_fov.x *= 1. - 0.09 * multiplier; - viewing_glass.viewport.w_h_fov.y *= 1. - 0.09 * multiplier; + tree_view.viewing_glass.viewport.w_h_fov.x *= 1. - 0.09 * multiplier; + tree_view.viewing_glass.viewport.w_h_fov.y *= 1. - 0.09 * multiplier; } } diff --git a/examples/minecraft.rs b/examples/minecraft.rs index c33f176..d040679 100644 --- a/examples/minecraft.rs +++ b/examples/minecraft.rs @@ -1,16 +1,10 @@ -#[cfg(feature = "bevy_wgpu")] -use shocovox_rs::octree::Octree; - #[cfg(feature = "bevy_wgpu")] use bevy::{prelude::*, window::WindowPlugin}; #[cfg(feature = "bevy_wgpu")] use shocovox_rs::octree::{ - raytracing::{ - bevy::create_viewing_glass, Ray, ShocoVoxRenderData, ShocoVoxRenderPlugin, - ShocoVoxViewingGlass, Viewport, - }, - Albedo, V3c, VoxelData, + raytracing::{OctreeGPUView, Ray, ShocoVoxRenderPlugin, Viewport}, + Albedo, Octree, V3c, VoxelData, }; #[cfg(feature = "bevy_wgpu")] @@ -85,9 +79,8 @@ fn setup(mut commands: Commands, images: ResMut>) { radius: tree.get_size() as f32 * 2.2, }); - let render_data = tree.create_bevy_view(); - let viewing_glass = create_viewing_glass( - &Viewport { + let output_texture = tree.create_new_gpu_view( + Viewport { origin: V3c { x: 0., y: 0., @@ -101,20 +94,20 @@ fn setup(mut commands: Commands, images: ResMut>) { w_h_fov: V3c::new(10., 10., 3.), }, DISPLAY_RESOLUTION, + &mut commands, images, ); + commands.spawn(SpriteBundle { sprite: Sprite { custom_size: Some(Vec2::new(1024., 768.)), ..default() }, - texture: viewing_glass.output_texture.clone(), + texture: output_texture, ..default() }); commands.spawn(Camera2dBundle::default()); commands.insert_resource(TreeResource { tree }); - commands.insert_resource(render_data); - commands.insert_resource(viewing_glass); commands.spawn(( PerfUiRoot::default(), @@ -144,35 +137,32 @@ struct DomePosition { } #[cfg(feature = "bevy_wgpu")] -fn rotate_camera( - angles_query: Query<&mut DomePosition>, - mut viewing_glass: ResMut, -) { +fn rotate_camera(angles_query: Query<&mut DomePosition>, mut tree_view: ResMut) { let (yaw, roll) = (angles_query.single().yaw, angles_query.single().roll); - let radius = angles_query.single().radius; - viewing_glass.viewport.origin = V3c::new( + + tree_view.viewing_glass.viewport.origin = V3c::new( radius / 2. + yaw.sin() * radius, radius + roll.sin() * radius * 2., radius / 2. + yaw.cos() * radius, ); - viewing_glass.viewport.direction = - (V3c::unit(radius / 2.) - viewing_glass.viewport.origin).normalized(); + tree_view.viewing_glass.viewport.direction = + (V3c::unit(radius / 2.) - tree_view.viewing_glass.viewport.origin).normalized(); } #[cfg(feature = "bevy_wgpu")] fn handle_zoom( keys: Res>, - mut viewing_glass: ResMut, - svx_data: Option>, + mut tree_view: ResMut, mut angles_query: Query<&mut DomePosition>, tree: Res>, ) { - if let Some(ref svx_data) = svx_data { - if svx_data.is_changed() { - println!("changed!! --> {:?}", svx_data.node_children[0..2].to_vec()); - } + if keys.pressed(KeyCode::Delete) { + tree_view.do_the_thing = true; } + // if tree_view.is_changed() { + // println!("changed!! --> {:?}", tree_view.read_back); + // } const ADDITION: f32 = 0.05; let angle_update_fn = |angle, delta| -> f32 { let new_angle = angle + delta; @@ -186,14 +176,17 @@ fn handle_zoom( // Render the current view with CPU let viewport_up_direction = V3c::new(0., 1., 0.); let viewport_right_direction = viewport_up_direction - .cross(viewing_glass.viewport.direction) + .cross(tree_view.viewing_glass.viewport.direction) .normalized(); - let pixel_width = viewing_glass.viewport.w_h_fov.x as f32 / DISPLAY_RESOLUTION[0] as f32; - let pixel_height = viewing_glass.viewport.w_h_fov.y as f32 / DISPLAY_RESOLUTION[1] as f32; - let viewport_bottom_left = viewing_glass.viewport.origin - + (viewing_glass.viewport.direction * viewing_glass.viewport.w_h_fov.z) - - (viewport_up_direction * (viewing_glass.viewport.w_h_fov.y / 2.)) - - (viewport_right_direction * (viewing_glass.viewport.w_h_fov.x / 2.)); + let pixel_width = + tree_view.viewing_glass.viewport.w_h_fov.x as f32 / DISPLAY_RESOLUTION[0] as f32; + let pixel_height = + tree_view.viewing_glass.viewport.w_h_fov.y as f32 / DISPLAY_RESOLUTION[1] as f32; + let viewport_bottom_left = tree_view.viewing_glass.viewport.origin + + (tree_view.viewing_glass.viewport.direction + * tree_view.viewing_glass.viewport.w_h_fov.z) + - (viewport_up_direction * (tree_view.viewing_glass.viewport.w_h_fov.y / 2.)) + - (viewport_right_direction * (tree_view.viewing_glass.viewport.w_h_fov.x / 2.)); // define light let diffuse_light_normal = V3c::new(0., -1., 1.).normalized(); @@ -211,8 +204,8 @@ fn handle_zoom( + viewport_right_direction * x as f32 * pixel_width + viewport_up_direction * y as f32 * pixel_height; let ray = Ray { - origin: viewing_glass.viewport.origin, - direction: (glass_point - viewing_glass.viewport.origin).normalized(), + origin: tree_view.viewing_glass.viewport.origin, + direction: (glass_point - tree_view.viewing_glass.viewport.origin).normalized(), }; use std::io::Write; @@ -269,15 +262,12 @@ fn handle_zoom( angles_query.single_mut().radius *= 1. + 0.02 * multiplier; } if keys.pressed(KeyCode::Home) { - viewing_glass.viewport.w_h_fov.x *= 1. + 0.09 * multiplier; - viewing_glass.viewport.w_h_fov.y *= 1. + 0.09 * multiplier; + tree_view.viewing_glass.viewport.w_h_fov.x *= 1. + 0.09 * multiplier; + tree_view.viewing_glass.viewport.w_h_fov.y *= 1. + 0.09 * multiplier; } if keys.pressed(KeyCode::End) { - viewing_glass.viewport.w_h_fov.x *= 1. - 0.09 * multiplier; - viewing_glass.viewport.w_h_fov.y *= 1. - 0.09 * multiplier; - } - if keys.pressed(KeyCode::Delete) { - svx_data.unwrap().do_the_thing = true; + tree_view.viewing_glass.viewport.w_h_fov.x *= 1. - 0.09 * multiplier; + tree_view.viewing_glass.viewport.w_h_fov.y *= 1. - 0.09 * multiplier; } } diff --git a/src/octree/convert/bytecode.rs b/src/octree/convert/bytecode.rs index 7210069..846b77c 100644 --- a/src/octree/convert/bytecode.rs +++ b/src/octree/convert/bytecode.rs @@ -409,6 +409,7 @@ where )?; let node_children = Vec::decode_bencode_object(list.next_object()?.unwrap())?; Ok(Self { + views: Vec::new(), auto_simplify, octree_size: root_size, nodes, diff --git a/src/octree/mod.rs b/src/octree/mod.rs index 0efe2d3..37dee73 100644 --- a/src/octree/mod.rs +++ b/src/octree/mod.rs @@ -82,6 +82,8 @@ where let root_node_key = nodes.push(NodeContent::Nothing); // The first element is the root Node assert!(root_node_key == 0); Ok(Self { + #[cfg(feature = "bevy_wgpu")] + views: Vec::new(), auto_simplify: true, octree_size: size, nodes, diff --git a/src/octree/node.rs b/src/octree/node.rs index 059aed2..ee4d90d 100644 --- a/src/octree/node.rs +++ b/src/octree/node.rs @@ -90,7 +90,7 @@ where ///#################################################################################### impl BrickData where - T: VoxelData + PartialEq + Clone + Copy + Default, + T: VoxelData + PartialEq + Clone + Default, { /// Provides occupancy information for the part of the brick corresponmding /// to the given octant based on the contents of the brick @@ -231,7 +231,7 @@ where if homogeneous_type.is_empty() { *self = BrickData::Empty; } else { - *self = BrickData::Solid(*homogeneous_type); + *self = BrickData::Solid(homogeneous_type.clone()); } true } else { diff --git a/src/octree/raytracing/bevy/data.rs b/src/octree/raytracing/bevy/data.rs index c4fba04..77a92ed 100644 --- a/src/octree/raytracing/bevy/data.rs +++ b/src/octree/raytracing/bevy/data.rs @@ -1,25 +1,152 @@ use crate::object_pool::empty_marker; use crate::{ octree::{ - raytracing::bevy::types::{ - OctreeMetaData, ShocoVoxRenderData, ShocoVoxRenderPipeline, Voxelement, + raytracing::bevy::{ + create_output_texture, + types::{ + OctreeGPUView, OctreeMetaData, ShocoVoxRenderData, ShocoVoxRenderPipeline, + ShocoVoxViewingGlass, Viewport, Voxelement, + }, }, types::{NodeChildrenArray, NodeContent}, - Albedo, BrickData, Octree, V3c, VoxelData, + BrickData, Octree, V3c, VoxelData, }, spatial::lut::BITMAP_MASK_FOR_OCTANT_LUT, }; use bevy::{ ecs::system::{Res, ResMut}, math::Vec4, + prelude::{Assets, Commands, Handle, Image}, render::renderer::RenderDevice, }; -use std::collections::HashMap; +use std::{ + collections::HashMap, + sync::{Arc, Mutex}, +}; + +use super::types::OctreeGPUDataHandler; impl Octree where T: Default + Clone + Copy + PartialEq + VoxelData, { + //############################################################################## + // ███████ █████████ ███████████ ███████████ ██████████ ██████████ + // ███░░░░░███ ███░░░░░███░█░░░███░░░█░░███░░░░░███ ░░███░░░░░█░░███░░░░░█ + // ███ ░░███ ███ ░░░ ░ ░███ ░ ░███ ░███ ░███ █ ░ ░███ █ ░ + // ░███ ░███░███ ░███ ░██████████ ░██████ ░██████ + // ░███ ░███░███ ░███ ░███░░░░░███ ░███░░█ ░███░░█ + // ░░███ ███ ░░███ ███ ░███ ░███ ░███ ░███ ░ █ ░███ ░ █ + // ░░░███████░ ░░█████████ █████ █████ █████ ██████████ ██████████ + // ░░░░░░░ ░░░░░░░░░ ░░░░░ ░░░░░ ░░░░░ ░░░░░░░░░░ ░░░░░░░░░░ + // █████████ ███████████ █████ █████ + // ███░░░░░███░░███░░░░░███░░███ ░░███ + // ███ ░░░ ░███ ░███ ░███ ░███ + // ░███ ░██████████ ░███ ░███ + // ░███ █████ ░███░░░░░░ ░███ ░███ + // ░░███ ░░███ ░███ ░███ ░███ + // ░░█████████ █████ ░░████████ + // ░░░░░░░░░ ░░░░░ ░░░░░░░░ + // █████ █████ █████ ██████████ █████ ███ █████ + // ░░███ ░░███ ░░███ ░░███░░░░░█░░███ ░███ ░░███ + // ░███ ░███ ░███ ░███ █ ░ ░███ ░███ ░███ + // ░███ ░███ ░███ ░██████ ░███ ░███ ░███ + // ░░███ ███ ░███ ░███░░█ ░░███ █████ ███ + // ░░░█████░ ░███ ░███ ░ █ ░░░█████░█████░ + // ░░███ █████ ██████████ ░░███ ░░███ + //############################################################################## + + /// Creates GPU compatible data renderable on the GPU from an octree + pub fn create_new_gpu_view( + &self, + viewport: Viewport, + resolution: [u32; 2], + commands: &mut Commands, + images: ResMut>, + ) -> Handle { + let mut gpu_data_handler = OctreeGPUDataHandler { + debug_gpu_interface: None, + readable_debug_gpu_interface: None, + render_data: ShocoVoxRenderData { + debug_gpu_interface: 0, + octree_meta: OctreeMetaData { + octree_size: self.octree_size, + voxel_brick_dim: DIM as u32, + ambient_light_color: V3c::new(1., 1., 1.), + ambient_light_position: V3c::new( + self.octree_size as f32, + self.octree_size as f32, + self.octree_size as f32, + ), + }, + metadata: Vec::with_capacity(self.nodes.len()), + node_children: Vec::with_capacity(self.nodes.len() * 8), + voxels: Vec::new(), + node_occupied_bits: Vec::with_capacity(self.nodes.len() * 2), + color_palette: Vec::new(), + }, + // victim_node_pointer: 0, + // victim_brick_pointer: 0, + map_to_color_index_in_palette: HashMap::new(), + map_to_node_index_in_nodes_buffer: HashMap::new(), + }; + + // Build up Nodes + for node_key in 0..self.nodes.len() { + if self.nodes.key_is_valid(node_key) { + gpu_data_handler.add_node_properties(&self, node_key); + } + } + + // Build up node content + for node_key in 0..self.nodes.len() { + if self.nodes.key_is_valid(node_key) { + gpu_data_handler.add_node_content(&self, node_key); + } + } + + let output_texture = create_output_texture(resolution, images); + commands.insert_resource(OctreeGPUView { + do_the_thing: false, + read_back: 0, + data_handler: Arc::new(Mutex::new(gpu_data_handler)), + viewing_glass: ShocoVoxViewingGlass { + output_texture: output_texture.clone(), + viewport: viewport, + }, + }); + output_texture + } +} + +//############################################################################## +// ███████ █████████ ███████████ ███████████ ██████████ ██████████ +// ███░░░░░███ ███░░░░░███░█░░░███░░░█░░███░░░░░███ ░░███░░░░░█░░███░░░░░█ +// ███ ░░███ ███ ░░░ ░ ░███ ░ ░███ ░███ ░███ █ ░ ░███ █ ░ +// ░███ ░███░███ ░███ ░██████████ ░██████ ░██████ +// ░███ ░███░███ ░███ ░███░░░░░███ ░███░░█ ░███░░█ +// ░░███ ███ ░░███ ███ ░███ ░███ ░███ ░███ ░ █ ░███ ░ █ +// ░░░███████░ ░░█████████ █████ █████ █████ ██████████ ██████████ +// ░░░░░░░ ░░░░░░░░░ ░░░░░ ░░░░░ ░░░░░ ░░░░░░░░░░ ░░░░░░░░░░ +// █████████ ███████████ █████ █████ +// ███░░░░░███░░███░░░░░███░░███ ░░███ +// ███ ░░░ ░███ ░███ ░███ ░███ +// ░███ ░██████████ ░███ ░███ +// ░███ █████ ░███░░░░░░ ░███ ░███ +// ░░███ ░░███ ░███ ░███ ░███ +// ░░█████████ █████ ░░████████ +// ░░░░░░░░░ ░░░░░ ░░░░░░░░ +// ██████████ █████████ ███████████ █████████ +// ░░███░░░░███ ███░░░░░███ ░█░░░███░░░█ ███░░░░░███ +// ░███ ░░███ ░███ ░███ ░ ░███ ░ ░███ ░███ +// ░███ ░███ ░███████████ ░███ ░███████████ +// ░███ ░███ ░███░░░░░███ ░███ ░███░░░░░███ +// ░███ ███ ░███ ░███ ░███ ░███ ░███ +// ██████████ █████ █████ █████ █████ █████ +// ░░░░░░░░░░ ░░░░░ ░░░░░ ░░░░░ ░░░░░ ░░░░░ +//############################################################################## + +impl OctreeGPUDataHandler { /// Updates metadata to set every element it might contain unused fn meta_set_element_unused(sized_node_meta: &mut u32) { *sized_node_meta = *sized_node_meta & 0x00FFFFFC; @@ -44,11 +171,13 @@ where /// * `sized_node_meta` - the bytes to update /// * `brick` - the brick to describe into the bytes /// * `brick_octant` - the octant to update in the bytes - fn meta_add_leaf_brick_structure( + fn meta_add_leaf_brick_structure( sized_node_meta: &mut u32, brick: &BrickData, brick_octant: usize, - ) { + ) where + T: Default + Clone + PartialEq + VoxelData, + { match brick { BrickData::Empty => {} // Child structure properties already set to NIL BrickData::Solid(_voxel) => { @@ -67,7 +196,12 @@ where /// Updates the given meta element value to store the leaf structure of the given node /// the given NodeContent reference is expected to be a leaf node - fn meta_set_leaf_structure(sized_node_meta: &mut u32, leaf: &NodeContent) { + fn meta_set_leaf_structure( + sized_node_meta: &mut u32, + leaf: &NodeContent, + ) where + T: Default + Copy + Clone + PartialEq + VoxelData, + { match leaf { NodeContent::UniformLeaf(brick) => { Self::meta_set_is_leaf(sized_node_meta, true); @@ -88,7 +222,10 @@ where } /// Creates the descriptor bytes for the given node - fn create_node_properties(node: &NodeContent) -> u32 { + fn create_node_properties(node: &NodeContent) -> u32 + where + T: Default + Copy + Clone + PartialEq + VoxelData, + { let mut meta = 0; Self::meta_set_element_unused(&mut meta); match node { @@ -103,43 +240,177 @@ where meta } + fn add_node_properties(&mut self, tree: &Octree, node_key: usize) + where + T: Default + Copy + Clone + PartialEq + VoxelData, + { + if tree.nodes.key_is_valid(node_key) { + self.map_to_node_index_in_nodes_buffer + .insert(node_key, self.render_data.metadata.len()); + self.render_data + .metadata + .push(Self::create_node_properties(tree.nodes.get(node_key))); + } else { + panic!("Trying to query invalid node key to set node metadata!"); + } + } + + fn add_node_content(&mut self, tree: &Octree, node_key: usize) + where + T: Default + Copy + Clone + PartialEq + VoxelData, + { + let occupied_bits = tree.stored_occupied_bits(node_key); + self.render_data.node_occupied_bits.extend_from_slice(&[ + (occupied_bits & 0x00000000FFFFFFFF) as u32, + ((occupied_bits & 0xFFFFFFFF00000000) >> 32) as u32, + ]); + match tree.nodes.get(node_key) { + NodeContent::UniformLeaf(brick) => { + debug_assert!( + matches!( + tree.node_children[node_key].content, + NodeChildrenArray::OccupancyBitmap(_) + ), + "Expected Uniform leaf to have OccupancyBitmap(_) instead of {:?}", + tree.node_children[node_key].content + ); + + let (brick_index, brick_added) = self.add_brick(brick); + + self.render_data.node_children.extend_from_slice(&[ + brick_index, + empty_marker(), + empty_marker(), + empty_marker(), + empty_marker(), + empty_marker(), + empty_marker(), + empty_marker(), + ]); + #[cfg(debug_assertions)] + { + if !brick_added { + // If no brick was added, the occupied bits should either be empty or full + if let NodeChildrenArray::OccupancyBitmap(occupied_bits) = + tree.node_children[node_key].content + { + debug_assert!(occupied_bits == 0 || occupied_bits == u64::MAX); + } + } + } + } + NodeContent::Leaf(bricks) => { + debug_assert!( + matches!( + tree.node_children[node_key].content, + NodeChildrenArray::OccupancyBitmap(_) + ), + "Expected Leaf to have OccupancyBitmaps(_) instead of {:?}", + tree.node_children[node_key].content + ); + + let mut children = vec![empty_marker(); 8]; + for octant in 0..8 { + let (brick_index, brick_added) = self.add_brick(&bricks[octant]); + + children[octant] = brick_index; + #[cfg(debug_assertions)] + { + if !brick_added { + // If no brick was added, the relevant occupied bits should either be empty or full + if let NodeChildrenArray::OccupancyBitmap(occupied_bits) = + tree.node_children[node_key].content + { + debug_assert!( + 0 == occupied_bits & BITMAP_MASK_FOR_OCTANT_LUT[octant] + || BITMAP_MASK_FOR_OCTANT_LUT[octant] + == occupied_bits & BITMAP_MASK_FOR_OCTANT_LUT[octant] + ); + } + } + } + } + self.render_data.node_children.extend_from_slice(&children); + } + NodeContent::Internal(_) => { + for c in 0..8 { + let child_index = &tree.node_children[node_key][c]; + if *child_index != empty_marker() { + debug_assert!(self + .map_to_node_index_in_nodes_buffer + .contains_key(&(*child_index as usize))); + self.render_data.node_children.push( + self.map_to_node_index_in_nodes_buffer[&(*child_index as usize)] as u32, + ); + } else { + self.render_data.node_children.push(*child_index); + } + } + } + NodeContent::Nothing => { + self.render_data + .node_children + .extend_from_slice(&[empty_marker(); 8]); + } + } + println!( + "bricks: {:?} <> nodes: {:?}", + self.render_data.voxels.len() / (DIM * DIM * DIM), + self.render_data.metadata.len() + ); + + // debug_assert_eq!( + // self.render_data.metadata.len() * 2, + // self.render_data.node_occupied_bits.len(), + // "Node occupancy bitmaps length({:?}) should match node count({:?})!", + // self.render_data.node_occupied_bits.len(), + // self.render_data.metadata.len(), + // ); + + // debug_assert_eq!( + // self.render_data.metadata.len() * 8, + // self.render_data.node_children.len(), + // "Node count({:?}) should match length of children buffer({:?})!", + // self.render_data.metadata.len(), + // self.render_data.node_children.len() + // ); + } + /// Loads a brick into the provided voxels vector and color palette /// * `brick` - The brick to upload /// * `voxels` - The destination buffer /// * `color_palette` - The used color palette /// * `map_to_color_index_in_palette` - Indexing helper for the color palette /// * `returns` - the identifier to set in @SizedNode and true if a new brick was aded to the voxels vector - fn add_brick_to_vec( - brick: &BrickData, - voxels: &mut Vec, - color_palette: &mut Vec, - map_to_color_index_in_palette: &mut HashMap, - ) -> (u32, bool) { + fn add_brick(&mut self, brick: &BrickData) -> (u32, bool) + where + T: Default + Clone + PartialEq + VoxelData, + { match brick { BrickData::Empty => (empty_marker(), false), BrickData::Solid(voxel) => { let albedo = voxel.albedo(); if let std::collections::hash_map::Entry::Vacant(e) = - map_to_color_index_in_palette.entry(albedo) + self.map_to_color_index_in_palette.entry(albedo) { - e.insert(color_palette.len()); - color_palette.push(Vec4::new( + e.insert(self.render_data.color_palette.len()); + self.render_data.color_palette.push(Vec4::new( albedo.r as f32 / 255., albedo.g as f32 / 255., albedo.b as f32 / 255., albedo.a as f32 / 255., )); } - (map_to_color_index_in_palette[&albedo] as u32, false) + (self.map_to_color_index_in_palette[&albedo] as u32, false) } BrickData::Parted(brick) => { - voxels.reserve(DIM * DIM * DIM); - let brick_index = voxels.len() / (DIM * DIM * DIM); + self.render_data.voxels.reserve(DIM * DIM * DIM); + let brick_index = self.render_data.voxels.len() / (DIM * DIM * DIM); debug_assert_eq!( - voxels.len() % (DIM * DIM * DIM), + self.render_data.voxels.len() % (DIM * DIM * DIM), 0, "Expected Voxel buffer length({:?}) to be divisble by {:?}", - voxels.len(), + self.render_data.voxels.len(), (DIM * DIM * DIM) ); for z in 0..DIM { @@ -147,19 +418,19 @@ where for x in 0..DIM { let albedo = brick[x][y][z].albedo(); if let std::collections::hash_map::Entry::Vacant(e) = - map_to_color_index_in_palette.entry(albedo) + self.map_to_color_index_in_palette.entry(albedo) { - e.insert(color_palette.len()); - color_palette.push(Vec4::new( + e.insert(self.render_data.color_palette.len()); + self.render_data.color_palette.push(Vec4::new( albedo.r as f32 / 255., albedo.g as f32 / 255., albedo.b as f32 / 255., albedo.a as f32 / 255., )); } - let albedo_index = map_to_color_index_in_palette[&albedo]; + let albedo_index = self.map_to_color_index_in_palette[&albedo]; - voxels.push(Voxelement { + self.render_data.voxels.push(Voxelement { albedo_index: albedo_index as u32, content: brick[x][y][z].user_data(), }); @@ -170,263 +441,130 @@ where } } } +} - /// Creates GPU compatible data renderable on the GPU from an octree - pub fn create_bevy_view(&self) -> ShocoVoxRenderData { - let mut nodes = Vec::with_capacity(self.nodes.len()); - let mut voxels = Vec::new(); - let mut node_occupied_bits = Vec::new(); - let mut color_palette = Vec::new(); - // Size of meta for one element is 2 Bytes, so array should be half + 1 for odd numbers - let mut node_children = Vec::with_capacity(self.nodes.len() * 8); - - // Build up Nodes - let mut map_to_node_index_in_nodes_buffer = HashMap::new(); - for i in 0..self.nodes.len() { - if self.nodes.key_is_valid(i) { - map_to_node_index_in_nodes_buffer.insert(i, nodes.len()); - nodes.push(Self::create_node_properties(self.nodes.get(i))); - } - } - - // Build up voxel content - let mut map_to_color_index_in_palette = HashMap::new(); - for i in 0..self.nodes.len() { - if !self.nodes.key_is_valid(i) { - continue; - } - let occupied_bits = self.stored_occupied_bits(i); - node_occupied_bits.extend_from_slice(&[ - (occupied_bits & 0x00000000FFFFFFFF) as u32, - ((occupied_bits & 0xFFFFFFFF00000000) >> 32) as u32, - ]); - match self.nodes.get(i) { - NodeContent::UniformLeaf(brick) => { - debug_assert!( - matches!( - self.node_children[i].content, - NodeChildrenArray::OccupancyBitmap(_) - ), - "Expected Uniform leaf to have OccupancyBitmap(_) instead of {:?}", - self.node_children[i].content - ); - - let (brick_index, brick_added) = Self::add_brick_to_vec( - brick, - &mut voxels, - &mut color_palette, - &mut map_to_color_index_in_palette, - ); - - node_children.extend_from_slice(&[ - brick_index, - empty_marker(), - empty_marker(), - empty_marker(), - empty_marker(), - empty_marker(), - empty_marker(), - empty_marker(), - ]); - #[cfg(debug_assertions)] - { - if !brick_added { - // If no brick was added, the occupied bits should either be empty or full - if let NodeChildrenArray::OccupancyBitmap(occupied_bits) = - self.node_children[i].content - { - debug_assert!(occupied_bits == 0 || occupied_bits == u64::MAX); - } - } - } - } - NodeContent::Leaf(bricks) => { - debug_assert!( - matches!( - self.node_children[i].content, - NodeChildrenArray::OccupancyBitmap(_) - ), - "Expected Leaf to have OccupancyBitmaps(_) instead of {:?}", - self.node_children[i].content - ); - - let mut children = vec![empty_marker(); 8]; - for octant in 0..8 { - let (brick_index, brick_added) = Self::add_brick_to_vec( - &bricks[octant], - &mut voxels, - &mut color_palette, - &mut map_to_color_index_in_palette, - ); - - children[octant] = brick_index; - #[cfg(debug_assertions)] - { - if !brick_added { - // If no brick was added, the relevant occupied bits should either be empty or full - if let NodeChildrenArray::OccupancyBitmap(occupied_bits) = - self.node_children[i].content - { - debug_assert!( - 0 == occupied_bits & BITMAP_MASK_FOR_OCTANT_LUT[octant] - || BITMAP_MASK_FOR_OCTANT_LUT[octant] - == occupied_bits - & BITMAP_MASK_FOR_OCTANT_LUT[octant] - ); - } - } - } - } - node_children.extend_from_slice(&children); - } - NodeContent::Internal(_) => { - for c in 0..8 { - let child_index = &self.node_children[i][c]; - if *child_index != empty_marker() { - debug_assert!(map_to_node_index_in_nodes_buffer - .contains_key(&(*child_index as usize))); - node_children.push( - map_to_node_index_in_nodes_buffer[&(*child_index as usize)] as u32, - ); - } else { - node_children.push(*child_index); - } - } - } - NodeContent::Nothing => { - node_children.extend_from_slice(&[empty_marker(); 8]); - } - } - } - - // +++ DEBUG +++ - println!( - "bricks: {:?} <> nodes: {:?}", - voxels.len() / (DIM * DIM * DIM), - nodes.len() - ); - nodes.push(0); // Additional element in node array for debug purposes - /*debug_assert_eq!( - nodes.len() * 2, - node_occupied_bits.len(), - "Node occupancy bitmaps length({:?}) should match node count({:?})!", - node_occupied_bits.len(), - nodes.len(), - ); - - debug_assert_eq!( - nodes.len() * 8, - node_children.len(), - "Node count({:?}) should match length of children buffer({:?})!", - nodes.len(), - node_children.len() - );*/ - // --- DEBUG --- - ShocoVoxRenderData { - do_the_thing: false, - octree_meta: OctreeMetaData { - octree_size: self.octree_size, - voxel_brick_dim: DIM as u32, - ambient_light_color: V3c::new(1., 1., 1.), - ambient_light_position: V3c::new( - self.octree_size as f32, - self.octree_size as f32, - self.octree_size as f32, - ), - }, - metadata: nodes, - node_children, - voxels, - node_occupied_bits, - color_palette, - } - } - - pub(crate) fn insert_elements_into_cache( - &self, - render_data: &mut ShocoVoxRenderData, - requested_nodes: Vec, - ) { - //TODO: find the first unused element, and overwrite it with the item +pub(crate) fn sync_with_main_world( + tree_view: Option>, + mut world: ResMut, +) { + if let Some(tree_view) = tree_view { + let mut tree_view_mainworld = world.get_resource_mut::().unwrap(); + tree_view_mainworld.read_back = tree_view.read_back; } } +//############################################################################## +// █████████ ███████████ █████ █████ +// ███░░░░░███░░███░░░░░███░░███ ░░███ +// ███ ░░░ ░███ ░███ ░███ ░███ +// ░███ ░██████████ ░███ ░███ +// ░███ █████ ░███░░░░░░ ░███ ░███ +// ░░███ ░░███ ░███ ░███ ░███ +// ░░█████████ █████ ░░████████ +// ░░░░░░░░░ ░░░░░ ░░░░░░░░ +// ███████████ ██████████ █████████ ██████████ +// ░░███░░░░░███ ░░███░░░░░█ ███░░░░░███ ░░███░░░░███ +// ░███ ░███ ░███ █ ░ ░███ ░███ ░███ ░░███ +// ░██████████ ░██████ ░███████████ ░███ ░███ +// ░███░░░░░███ ░███░░█ ░███░░░░░███ ░███ ░███ +// ░███ ░███ ░███ ░ █ ░███ ░███ ░███ ███ +// █████ █████ ██████████ █████ █████ ██████████ +// ░░░░░ ░░░░░ ░░░░░░░░░░ ░░░░░ ░░░░░ ░░░░░░░░░░ +//############################################################################## pub(crate) fn handle_gpu_readback( render_device: Res, - svx_data: Option>, - svx_pipeline: Option>, + tree_gpu_view: Option>, ) { - // // Data updates triggered by debug interface - // if let Some(mut svx_data) = svx_data { - // let svx_pipeline = svx_pipeline.unwrap(); - // if svx_data.do_the_thing { - // // GPU buffer read - // // https://docs.rs/bevy/latest/src/gpu_readback/gpu_readback.rs.html - // let buffer_slice = svx_pipeline - // .readable_nodes_buffer - // .as_ref() - // .unwrap() - // .slice(..); - // let (s, r) = crossbeam::channel::unbounded::<()>(); - // buffer_slice.map_async( - // bevy::render::render_resource::MapMode::Read, - // move |d| match d { - // Ok(_) => s.send(()).expect("Failed to send map update"), - // Err(err) => println!("Something's wrong: {err}"), - // }, - // ); - - // render_device - // .poll(bevy::render::render_resource::Maintain::wait()) - // .panic_on_timeout(); - - // r.recv().expect("Failed to receive the map_async message"); - // { - // let buffer_view = buffer_slice.get_mapped_range(); + // Data updates triggered by debug interface + if let Some(mut tree_gpu_view) = tree_gpu_view { + if tree_gpu_view.do_the_thing { + let received_data; + { + let data_handler = tree_gpu_view.data_handler.lock().unwrap(); + // GPU buffer read + // https://docs.rs/bevy/latest/src/gpu_readback/gpu_readback.rs.html + let buffer_slice = data_handler + .readable_debug_gpu_interface + .as_ref() + .unwrap() + .slice(..); + let (s, r) = crossbeam::channel::unbounded::<()>(); + buffer_slice.map_async(bevy::render::render_resource::MapMode::Read, move |d| { + match d { + Ok(_) => s.send(()).expect("Failed to send map update"), + Err(err) => println!("Something's wrong: {err}"), + } + }); - // let data = buffer_view - // .chunks(std::mem::size_of::()) - // .map(|chunk| u32::from_ne_bytes(chunk.try_into().expect("should be a u32"))) - // .collect::>(); - // println!("data: {}", data[0]); - // } + render_device + .poll(bevy::render::render_resource::Maintain::wait()) + .panic_on_timeout(); - // svx_pipeline.readable_nodes_buffer.as_ref().unwrap().unmap(); + r.recv().expect("Failed to receive the map_async message"); + { + let buffer_view = buffer_slice.get_mapped_range(); - // svx_data.do_the_thing = false; - // } - // } + let data = buffer_view + .chunks(std::mem::size_of::()) + .map(|chunk| u32::from_ne_bytes(chunk.try_into().expect("should be a u32"))) + .collect::>(); + received_data = data[0]; + } + data_handler + .readable_debug_gpu_interface + .as_ref() + .unwrap() + .unmap(); + } + tree_gpu_view.read_back = received_data; + tree_gpu_view.do_the_thing = false; + } + } } -pub(crate) fn sync_with_main_world(// svx_data: Option>, - // svx_pipeline: Option>, - // mut world: ResMut, -) { - // let mut render_data_mainworld = world.get_resource_mut::().unwrap(); -} +//############################################################################## +// █████████ ███████████ █████ █████ +// ███░░░░░███░░███░░░░░███░░███ ░░███ +// ███ ░░░ ░███ ░███ ░███ ░███ +// ░███ ░██████████ ░███ ░███ +// ░███ █████ ░███░░░░░░ ░███ ░███ +// ░░███ ░░███ ░███ ░███ ░███ +// ░░█████████ █████ ░░████████ +// ░░░░░░░░░ ░░░░░ ░░░░░░░░ -pub(crate) fn handle_cache( - svx_data: Option>, +// █████ ███ █████ ███████████ █████ ███████████ ██████████ +// ░░███ ░███ ░░███ ░░███░░░░░███ ░░███ ░█░░░███░░░█░░███░░░░░█ +// ░███ ░███ ░███ ░███ ░███ ░███ ░ ░███ ░ ░███ █ ░ +// ░███ ░███ ░███ ░██████████ ░███ ░███ ░██████ +// ░░███ █████ ███ ░███░░░░░███ ░███ ░███ ░███░░█ +// ░░░█████░█████░ ░███ ░███ ░███ ░███ ░███ ░ █ +// ░░███ ░░███ █████ █████ █████ █████ ██████████ +// ░░░ ░░░ ░░░░░ ░░░░░ ░░░░░ ░░░░░ ░░░░░░░░░░ +//############################################################################## +pub(crate) fn write_to_gpu( + tree_gpu_view: Option>, svx_pipeline: Option>, ) { //TODO: Document that all components are lost during extract transition // Data updates triggered by debug interface - if let Some(mut svx_data) = svx_data { + if let Some(tree_gpu_view) = tree_gpu_view { let svx_pipeline = svx_pipeline.unwrap(); let render_queue = &svx_pipeline.render_queue.0; - if svx_data.do_the_thing { + if tree_gpu_view.do_the_thing { // GPU buffer Write - let data: u32 = 1; use bevy::render::render_resource::encase::StorageBuffer; - let mut data_buffer = StorageBuffer::new(Vec::::new()); - data_buffer.write(&data).unwrap(); + let data_buffer = StorageBuffer::new(vec![0, 0, 0, 1]); render_queue.write_buffer( - svx_pipeline.metadata_buffer.as_ref().unwrap(), - ((svx_data.metadata.len() - 1) * std::mem::size_of::()) as u64, + tree_gpu_view + .data_handler + .lock() + .unwrap() + .debug_gpu_interface + .as_ref() + .unwrap(), + 0, &data_buffer.into_inner(), ); - - svx_data.do_the_thing = false; } } } diff --git a/src/octree/raytracing/bevy/mod.rs b/src/octree/raytracing/bevy/mod.rs index d0ef1c6..18ba83c 100644 --- a/src/octree/raytracing/bevy/mod.rs +++ b/src/octree/raytracing/bevy/mod.rs @@ -3,13 +3,13 @@ mod pipeline; pub mod types; pub use crate::octree::raytracing::bevy::types::{ - ShocoVoxRenderPlugin, ShocoVoxViewingGlass, Viewport, + OctreeGPUView, ShocoVoxRenderPlugin, ShocoVoxViewingGlass, Viewport, }; use crate::octree::raytracing::bevy::{ - data::{handle_cache, handle_gpu_readback, sync_with_main_world}, + data::{handle_gpu_readback, sync_with_main_world, write_to_gpu}, pipeline::prepare_bind_groups, - types::{ShocoVoxLabel, ShocoVoxRenderData, ShocoVoxRenderNode, ShocoVoxRenderPipeline}, + types::{ShocoVoxLabel, ShocoVoxRenderNode, ShocoVoxRenderPipeline}, }; use bevy::{ @@ -27,18 +27,7 @@ use bevy::{ }, }; -pub fn create_viewing_glass( - viewport: &Viewport, - resolution: [u32; 2], - images: ResMut>, -) -> ShocoVoxViewingGlass { - ShocoVoxViewingGlass { - output_texture: create_ouput_texture(resolution, images), - viewport: *viewport, - } -} - -pub(crate) fn create_ouput_texture( +pub(crate) fn create_output_texture( resolution: [u32; 2], mut images: ResMut>, ) -> Handle { @@ -60,14 +49,13 @@ pub(crate) fn create_ouput_texture( impl Plugin for ShocoVoxRenderPlugin { fn build(&self, app: &mut App) { - app.add_plugins(ExtractResourcePlugin::::default()); - app.add_plugins(ExtractResourcePlugin::::default()); + app.add_plugins(ExtractResourcePlugin::::default()); let render_app = app.sub_app_mut(RenderApp); render_app.add_systems(ExtractSchedule, sync_with_main_world); render_app.add_systems( Render, ( - handle_cache.in_set(RenderSet::PrepareResources), + write_to_gpu.in_set(RenderSet::PrepareResources), prepare_bind_groups.in_set(RenderSet::PrepareBindGroups), handle_gpu_readback.in_set(RenderSet::Cleanup), ), diff --git a/src/octree/raytracing/bevy/pipeline.rs b/src/octree/raytracing/bevy/pipeline.rs index 9dde0e7..98204f7 100644 --- a/src/octree/raytracing/bevy/pipeline.rs +++ b/src/octree/raytracing/bevy/pipeline.rs @@ -1,5 +1,6 @@ use crate::octree::raytracing::bevy::types::{ - ShocoVoxRenderData, ShocoVoxRenderNode, ShocoVoxRenderPipeline, ShocoVoxViewingGlass, + OctreeGPUView, ShocoVoxRenderData, ShocoVoxRenderNode, ShocoVoxRenderPipeline, + ShocoVoxViewingGlass, }; use bevy::{ @@ -13,7 +14,7 @@ use bevy::{ render_graph::{self}, render_resource::{ encase::{StorageBuffer, UniformBuffer}, - AsBindGroup, BufferInitDescriptor, BufferUsages, CachedPipelineState, + AsBindGroup, BufferDescriptor, BufferInitDescriptor, BufferUsages, CachedPipelineState, ComputePassDescriptor, ComputePipelineDescriptor, PipelineCache, }, renderer::{RenderContext, RenderDevice, RenderQueue}, @@ -22,6 +23,8 @@ use bevy::{ }; use std::borrow::Cow; +use super::types::OctreeGPUDataHandler; + impl FromWorld for ShocoVoxRenderPipeline { fn from_world(world: &mut World) -> Self { let render_device = world.resource::(); @@ -45,7 +48,7 @@ impl FromWorld for ShocoVoxRenderPipeline { }); ShocoVoxRenderPipeline { - victim_pointer: 0, + tree_data_handler: None, render_queue: world.resource::().clone(), octree_meta_buffer: None, metadata_buffer: None, @@ -64,18 +67,28 @@ impl FromWorld for ShocoVoxRenderPipeline { } } +//############################################################################## +// ███████████ █████ █████ ██████ █████ +// ░░███░░░░░███ ░░███ ░░███ ░░██████ ░░███ +// ░███ ░███ ░███ ░███ ░███░███ ░███ +// ░██████████ ░███ ░███ ░███░░███░███ +// ░███░░░░░███ ░███ ░███ ░███ ░░██████ +// ░███ ░███ ░███ ░███ ░███ ░░█████ +// █████ █████ ░░████████ █████ ░░█████ +// ░░░░░ ░░░░░ ░░░░░░░░ ░░░░░ ░░░░░ +//############################################################################## const WORKGROUP_SIZE: u32 = 8; impl render_graph::Node for ShocoVoxRenderNode { fn update(&mut self, world: &mut World) { { - let render_data = world.get_resource::(); let svx_pipeline = world.resource::(); let pipeline_cache = world.resource::(); + let tree_gpu_view = world.get_resource::(); if !self.ready { if let CachedPipelineState::Ok(_) = pipeline_cache.get_compute_pipeline_state(svx_pipeline.update_pipeline) { - self.ready = render_data.is_some(); + self.ready = tree_gpu_view.is_some(); } } } @@ -119,38 +132,172 @@ impl render_graph::Node for ShocoVoxRenderNode { svx_pipeline.readable_metadata_buffer.as_ref().unwrap(), 0, std::mem::size_of::() as u64, + ); + // +++ DEBUG +++ + let tree_gpu_view = world.resource::(); + let data_handler = tree_gpu_view.data_handler.lock().unwrap(); + command_encoder.copy_buffer_to_buffer( + data_handler.debug_gpu_interface.as_ref().unwrap(), + 0, + data_handler.readable_debug_gpu_interface.as_ref().unwrap(), + 0, + std::mem::size_of::() as u64, ) + // --- DEBUG --- } Ok(()) } } +//############################################################################## +// █████████ ███████████ ██████████ █████████ ███████████ ██████████ +// ███░░░░░███░░███░░░░░███ ░░███░░░░░█ ███░░░░░███ ░█░░░███░░░█░░███░░░░░█ +// ███ ░░░ ░███ ░███ ░███ █ ░ ░███ ░███ ░ ░███ ░ ░███ █ ░ +// ░███ ░██████████ ░██████ ░███████████ ░███ ░██████ +// ░███ ░███░░░░░███ ░███░░█ ░███░░░░░███ ░███ ░███░░█ +// ░░███ ███ ░███ ░███ ░███ ░ █ ░███ ░███ ░███ ░███ ░ █ +// ░░█████████ █████ █████ ██████████ █████ █████ █████ ██████████ +// ░░░░░░░░░ ░░░░░ ░░░░░ ░░░░░░░░░░ ░░░░░ ░░░░░ ░░░░░ ░░░░░░░░░░ +// ███████████ █████ ██████ █████ ██████████ +// ░░███░░░░░███░░███ ░░██████ ░░███ ░░███░░░░███ +// ░███ ░███ ░███ ░███░███ ░███ ░███ ░░███ +// ░██████████ ░███ ░███░░███░███ ░███ ░███ +// ░███░░░░░███ ░███ ░███ ░░██████ ░███ ░███ +// ░███ ░███ ░███ ░███ ░░█████ ░███ ███ +// ███████████ █████ █████ ░░█████ ██████████ +// ░░░░░░░░░░░ ░░░░░ ░░░░░ ░░░░░ ░░░░░░░░░░ +// █████████ ███████████ ███████ █████ █████ ███████████ █████████ +// ███░░░░░███░░███░░░░░███ ███░░░░░███ ░░███ ░░███ ░░███░░░░░███ ███░░░░░███ +// ███ ░░░ ░███ ░███ ███ ░░███ ░███ ░███ ░███ ░███░███ ░░░ +// ░███ ░██████████ ░███ ░███ ░███ ░███ ░██████████ ░░█████████ +// ░███ █████ ░███░░░░░███ ░███ ░███ ░███ ░███ ░███░░░░░░ ░░░░░░░░███ +// ░░███ ░░███ ░███ ░███ ░░███ ███ ░███ ░███ ░███ ███ ░███ +// ░░█████████ █████ █████ ░░░███████░ ░░████████ █████ ░░█████████ +// ░░░░░░░░░ ░░░░░ ░░░░░ ░░░░░░░ ░░░░░░░░ ░░░░░ ░░░░░░░░░ +//############################################################################## pub(crate) fn prepare_bind_groups( gpu_images: Res>, fallback_image: Res, render_device: Res, mut pipeline: ResMut, - octree_viewing_glass: Res, - render_data: Res, + tree_gpu_view: ResMut, ) { - let bind_group = octree_viewing_glass - .as_bind_group( - &pipeline.viewing_glass_bind_group_layout, - &render_device, - &gpu_images, - &fallback_image, - ) - .ok() - .unwrap(); - pipeline.viewing_glass_bind_group = Some(bind_group.bind_group); + pipeline.viewing_glass_bind_group = Some( + tree_gpu_view + .viewing_glass + .as_bind_group( + &pipeline.viewing_glass_bind_group_layout, + &render_device, + &gpu_images, + &fallback_image, + ) + .ok() + .unwrap() + .bind_group, + ); if pipeline.update_tree { + let mut data_handler = tree_gpu_view.data_handler.lock().unwrap(); + + // Create the staging buffer helping in reading data from the GPU + //############################################################################## + // █████████ ███████████ █████ █████ + // ███░░░░░███░░███░░░░░███░░███ ░░███ + // ███ ░░░ ░███ ░███ ░███ ░███ + // ░███ ░██████████ ░███ ░███ + // ░███ █████ ░███░░░░░░ ░███ ░███ + // ░░███ ░░███ ░███ ░███ ░███ + // ░░█████████ █████ ░░████████ + // ░░░░░░░░░ ░░░░░ ░░░░░░░░ + // ███████████ ██████████ █████████ ██████████ + // ░░███░░░░░███ ░░███░░░░░█ ███░░░░░███ ░░███░░░░███ + // ░███ ░███ ░███ █ ░ ░███ ░███ ░███ ░░███ + // ░██████████ ░██████ ░███████████ ░███ ░███ + // ░███░░░░░███ ░███░░█ ░███░░░░░███ ░███ ░███ + // ░███ ░███ ░███ ░ █ ░███ ░███ ░███ ███ + // █████ █████ ██████████ █████ █████ ██████████ + // ░░░░░ ░░░░░ ░░░░░░░░░░ ░░░░░ ░░░░░ ░░░░░░░░░░ + // ███████████ █████ █████ ███████████ ███████████ ██████████ ███████████ + // ░░███░░░░░███░░███ ░░███ ░░███░░░░░░█░░███░░░░░░█░░███░░░░░█░░███░░░░░███ + // ░███ ░███ ░███ ░███ ░███ █ ░ ░███ █ ░ ░███ █ ░ ░███ ░███ + // ░██████████ ░███ ░███ ░███████ ░███████ ░██████ ░██████████ + // ░███░░░░░███ ░███ ░███ ░███░░░█ ░███░░░█ ░███░░█ ░███░░░░░███ + // ░███ ░███ ░███ ░███ ░███ ░ ░███ ░ ░███ ░ █ ░███ ░███ + // ███████████ ░░████████ █████ █████ ██████████ █████ █████ + // ░░░░░░░░░░░ ░░░░░░░░ ░░░░░ ░░░░░ ░░░░░░░░░░ ░░░░░ ░░░░░ + //############################################################################## + + if pipeline.readable_metadata_buffer.is_none() { + pipeline.readable_metadata_buffer = + Some(render_device.create_buffer(&BufferDescriptor { + mapped_at_creation: false, + size: (data_handler.render_data.metadata.len() * 4) as u64, + label: Some("Octree Node metadata Buffer"), + usage: BufferUsages::COPY_DST | BufferUsages::MAP_READ, + })); + } + + // +++ DEBUG +++ + if data_handler.readable_debug_gpu_interface.is_none() { + data_handler.readable_debug_gpu_interface = + Some(render_device.create_buffer(&BufferDescriptor { + mapped_at_creation: false, + size: 4, + label: Some("Octree Debug interface Buffer"), + usage: BufferUsages::COPY_DST | BufferUsages::MAP_READ, + })); + } + // --- DEBUG --- + + //############################################################################## + // ███████████ ██████████ ██████ █████ ██████████ ██████████ ███████████ + // ░░███░░░░░███ ░░███░░░░░█░░██████ ░░███ ░░███░░░░███ ░░███░░░░░█░░███░░░░░███ + // ░███ ░███ ░███ █ ░ ░███░███ ░███ ░███ ░░███ ░███ █ ░ ░███ ░███ + // ░██████████ ░██████ ░███░░███░███ ░███ ░███ ░██████ ░██████████ + // ░███░░░░░███ ░███░░█ ░███ ░░██████ ░███ ░███ ░███░░█ ░███░░░░░███ + // ░███ ░███ ░███ ░ █ ░███ ░░█████ ░███ ███ ░███ ░ █ ░███ ░███ + // █████ █████ ██████████ █████ ░░█████ ██████████ ██████████ █████ █████ + // ░░░░░ ░░░░░ ░░░░░░░░░░ ░░░░░ ░░░░░ ░░░░░░░░░░ ░░░░░░░░░░ ░░░░░ ░░░░░ + // ██████████ █████████ ███████████ █████████ + // ░░███░░░░███ ███░░░░░███ ░█░░░███░░░█ ███░░░░░███ + // ░███ ░░███ ░███ ░███ ░ ░███ ░ ░███ ░███ + // ░███ ░███ ░███████████ ░███ ░███████████ + // ░███ ░███ ░███░░░░░███ ░███ ░███░░░░░███ + // ░███ ███ ░███ ░███ ░███ ░███ ░███ + // ██████████ █████ █████ █████ █████ █████ + // ░░░░░░░░░░ ░░░░░ ░░░░░ ░░░░░ ░░░░░ ░░░░░ + // ███████████ █████ █████ ███████████ ███████████ ██████████ ███████████ + // ░░███░░░░░███░░███ ░░███ ░░███░░░░░░█░░███░░░░░░█░░███░░░░░█░░███░░░░░███ + // ░███ ░███ ░███ ░███ ░███ █ ░ ░███ █ ░ ░███ █ ░ ░███ ░███ + // ░██████████ ░███ ░███ ░███████ ░███████ ░██████ ░██████████ + // ░███░░░░░███ ░███ ░███ ░███░░░█ ░███░░░█ ░███░░█ ░███░░░░░███ + // ░███ ░███ ░███ ░███ ░███ ░ ░███ ░ ░███ ░ █ ░███ ░███ + // ███████████ ░░████████ █████ █████ ██████████ █████ █████ + // ░░░░░░░░░░░ ░░░░░░░░ ░░░░░ ░░░░░ ░░░░░░░░░░ ░░░░░ ░░░░░ + //############################################################################## + //================================================================= // Implementation with WGPU //================================================================= // Upload data to buffers + // +++ DEBUG +++ + let buffer = UniformBuffer::new(vec![0, 0, 0, 0]); + if let Some(debug_buffer) = &data_handler.debug_gpu_interface { + pipeline + .render_queue + .write_buffer(debug_buffer, 0, &buffer.into_inner()) + } else { + let debug_buffer = render_device.create_buffer_with_data(&BufferInitDescriptor { + label: Some("Octree Debug Buffer"), + contents: &buffer.into_inner(), + usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC, + }); + data_handler.debug_gpu_interface = Some(debug_buffer); + } + // --- DEBUG --- + let mut buffer = UniformBuffer::new(Vec::::new()); - buffer.write(&render_data.octree_meta).unwrap(); + buffer.write(&data_handler.render_data.octree_meta).unwrap(); if let Some(metadata_buffer) = &pipeline.octree_meta_buffer { pipeline .render_queue @@ -165,7 +312,7 @@ pub(crate) fn prepare_bind_groups( } let mut buffer = StorageBuffer::new(Vec::::new()); - buffer.write(&render_data.metadata).unwrap(); + buffer.write(&data_handler.render_data.metadata).unwrap(); if let Some(nodes_buffer) = &pipeline.metadata_buffer { pipeline .render_queue @@ -180,7 +327,9 @@ pub(crate) fn prepare_bind_groups( } let mut buffer = StorageBuffer::new(Vec::::new()); - buffer.write(&render_data.node_children).unwrap(); + buffer + .write(&data_handler.render_data.node_children) + .unwrap(); if let Some(children_buffer) = &pipeline.node_children_buffer { pipeline .render_queue @@ -195,7 +344,9 @@ pub(crate) fn prepare_bind_groups( } let mut buffer = StorageBuffer::new(Vec::::new()); - buffer.write(&render_data.node_occupied_bits).unwrap(); + buffer + .write(&data_handler.render_data.node_occupied_bits) + .unwrap(); if let Some(ocbits_buffer) = &pipeline.node_ocbits_buffer { pipeline .render_queue @@ -210,7 +361,7 @@ pub(crate) fn prepare_bind_groups( } let mut buffer = StorageBuffer::new(Vec::::new()); - buffer.write(&render_data.voxels).unwrap(); + buffer.write(&data_handler.render_data.voxels).unwrap(); if let Some(voxels_buffer) = &pipeline.voxels_buffer { pipeline .render_queue @@ -225,7 +376,9 @@ pub(crate) fn prepare_bind_groups( } let mut buffer = StorageBuffer::new(Vec::::new()); - buffer.write(&render_data.color_palette).unwrap(); + buffer + .write(&data_handler.render_data.color_palette) + .unwrap(); if let Some(color_palette_buffer) = &pipeline.color_palette_buffer { pipeline .render_queue @@ -240,6 +393,24 @@ pub(crate) fn prepare_bind_groups( pipeline.color_palette_buffer = Some(color_palette_buffer); } + //############################################################################## + // ███████████ █████ ██████ █████ ██████████ + // ░░███░░░░░███░░███ ░░██████ ░░███ ░░███░░░░███ + // ░███ ░███ ░███ ░███░███ ░███ ░███ ░░███ + // ░██████████ ░███ ░███░░███░███ ░███ ░███ + // ░███░░░░░███ ░███ ░███ ░░██████ ░███ ░███ + // ░███ ░███ ░███ ░███ ░░█████ ░███ ███ + // ███████████ █████ █████ ░░█████ ██████████ + // ░░░░░░░░░░░ ░░░░░ ░░░░░ ░░░░░ ░░░░░░░░░░ + // █████████ ███████████ ███████ █████ █████ ███████████ + // ███░░░░░███░░███░░░░░███ ███░░░░░███ ░░███ ░░███ ░░███░░░░░███ + // ███ ░░░ ░███ ░███ ███ ░░███ ░███ ░███ ░███ ░███ + // ░███ ░██████████ ░███ ░███ ░███ ░███ ░██████████ + // ░███ █████ ░███░░░░░███ ░███ ░███ ░███ ░███ ░███░░░░░░ + // ░░███ ░░███ ░███ ░███ ░░███ ███ ░███ ░███ ░███ + // ░░█████████ █████ █████ ░░░███████░ ░░████████ █████ + // ░░░░░░░░░ ░░░░░ ░░░░░ ░░░░░░░ ░░░░░░░░ ░░░░░ + //############################################################################## // Create bind group let tree_bind_group = render_device.create_bind_group( ShocoVoxRenderData::label(), @@ -289,6 +460,16 @@ pub(crate) fn prepare_bind_groups( .unwrap() .as_entire_binding(), }, + // +++ DEBUG +++ + bevy::render::render_resource::BindGroupEntry { + binding: 6, + resource: data_handler + .debug_gpu_interface + .as_ref() + .unwrap() + .as_entire_binding(), + }, + // --- DEBUG --- ], ); pipeline.tree_bind_group = Some(tree_bind_group); @@ -324,24 +505,6 @@ pub(crate) fn prepare_bind_groups( // } // pipeline.tree_bind_group = Some(tree_bind_group.bind_group); - - // Create the staging buffer helping in reading data from the GPU - let mut buffer = StorageBuffer::new(Vec::::new()); - buffer.write(&render_data.metadata).unwrap(); - if let Some(readable_nodes_buffer) = &pipeline.readable_metadata_buffer { - pipeline - .render_queue - .write_buffer(readable_nodes_buffer, 0, &buffer.into_inner()) - } else { - let readable_nodes_buffer = - render_device.create_buffer_with_data(&BufferInitDescriptor { - label: Some("Octree Cache Bytes Buffer"), - contents: &buffer.into_inner(), - usage: BufferUsages::COPY_DST | BufferUsages::MAP_READ, - }); - pipeline.readable_metadata_buffer = Some(readable_nodes_buffer); - } - pipeline.update_tree = false; } } diff --git a/src/octree/raytracing/bevy/types.rs b/src/octree/raytracing/bevy/types.rs index 007d5d7..5e0d057 100644 --- a/src/octree/raytracing/bevy/types.rs +++ b/src/octree/raytracing/bevy/types.rs @@ -1,4 +1,4 @@ -use crate::octree::V3cf32; +use crate::octree::{Albedo, V3cf32}; use bevy::{ asset::Handle, ecs::system::Resource, @@ -14,6 +14,10 @@ use bevy::{ renderer::RenderQueue, }, }; +use std::{ + collections::HashMap, + sync::{Arc, Mutex}, +}; #[derive(Clone, ShaderType)] pub(crate) struct Voxelement { @@ -41,6 +45,38 @@ pub struct ShocoVoxRenderPlugin { } #[derive(Resource, Clone, AsBindGroup, TypePath, ExtractResource)] +#[type_path = "shocovox::gpu::OctreeGPUView"] +pub struct OctreeGPUView { + // +++ DEBUG +++ + pub do_the_thing: bool, + pub read_back: u32, + // --- DEBUG --- + pub viewing_glass: ShocoVoxViewingGlass, + pub(crate) data_handler: Arc>, +} + +#[derive(Resource, Clone, AsBindGroup, TypePath, ExtractResource)] +#[type_path = "shocovox::gpu::OctreeGPUDataHandler"] +pub struct OctreeGPUDataHandler { + pub(crate) render_data: ShocoVoxRenderData, + // pub(crate) victim_node_pointer: u32, + // pub(crate) victim_brick_pointer: u32, + pub(crate) map_to_node_index_in_nodes_buffer: HashMap, + pub(crate) map_to_color_index_in_palette: HashMap, + pub(crate) debug_gpu_interface: Option, + pub(crate) readable_debug_gpu_interface: Option, + //TODO: Maybe this? + // Buffers for the RenderData + // pub(crate) octree_meta_buffer: Buffer, + // pub(crate) metadata_buffer: Buffer, + // pub(crate) node_children_buffer: Buffer, + // pub(crate) node_ocbits_buffer: Buffer, + // pub(crate) voxels_buffer: Buffer, + // pub(crate) color_palette_buffer: Buffer, + // pub(crate) readable_metadata_buffer: Buffer, +} + +#[derive(Clone, AsBindGroup, TypePath)] #[type_path = "shocovox::gpu::ShocoVoxViewingGlass"] pub struct ShocoVoxViewingGlass { #[storage_texture(0, image_format = Rgba8Unorm, access = ReadWrite)] @@ -50,11 +86,13 @@ pub struct ShocoVoxViewingGlass { pub viewport: Viewport, } -#[derive(Resource, Clone, AsBindGroup, TypePath, ExtractResource)] +#[derive(Clone, AsBindGroup, TypePath)] #[type_path = "shocovox::gpu::ShocoVoxRenderData"] pub struct ShocoVoxRenderData { - pub do_the_thing: bool, //STRICTLY FOR DEBUG REASONS - + // +++ DEBUG +++ + #[storage(6, visibility(compute))] + pub(crate) debug_gpu_interface: u32, + // --- DEBUG --- /// Contains the properties of the Octree #[uniform(0, visibility(compute))] pub(crate) octree_meta: OctreeMetaData, @@ -125,9 +163,6 @@ pub struct ShocoVoxRenderData { pub(crate) struct ShocoVoxRenderPipeline { pub update_tree: bool, - // The candidates for deletion inside nodes array on page deletion - pub(crate) victim_pointer: u32, - pub(crate) render_queue: RenderQueue, pub(crate) update_pipeline: CachedComputePipelineId, @@ -139,6 +174,7 @@ pub(crate) struct ShocoVoxRenderPipeline { pub(crate) node_ocbits_buffer: Option, pub(crate) voxels_buffer: Option, pub(crate) color_palette_buffer: Option, + pub(crate) tree_data_handler: Option, pub(crate) viewing_glass_bind_group_layout: BindGroupLayout, pub(crate) render_data_bind_group_layout: BindGroupLayout, diff --git a/src/octree/raytracing/mod.rs b/src/octree/raytracing/mod.rs index 7881c67..8615da6 100644 --- a/src/octree/raytracing/mod.rs +++ b/src/octree/raytracing/mod.rs @@ -7,4 +7,6 @@ pub mod bevy; pub use crate::spatial::raytracing::Ray; #[cfg(feature = "bevy_wgpu")] -pub use bevy::types::{ShocoVoxRenderData, ShocoVoxRenderPlugin, ShocoVoxViewingGlass, Viewport}; +pub use bevy::types::{ + OctreeGPUView, ShocoVoxRenderData, ShocoVoxRenderPlugin, ShocoVoxViewingGlass, Viewport, +}; diff --git a/src/octree/types.rs b/src/octree/types.rs index 2b05283..6cbfbf8 100644 --- a/src/octree/types.rs +++ b/src/octree/types.rs @@ -4,6 +4,9 @@ use std::error::Error; #[cfg(feature = "serialization")] use serde::{Deserialize, Serialize}; +#[cfg(feature = "bevy_wgpu")] +use {crate::octree::raytracing::bevy::types::OctreeGPUDataHandler, std::sync::Arc}; + #[derive(Debug, Clone, PartialEq)] #[cfg_attr(feature = "serialization", derive(Serialize, Deserialize))] pub(crate) enum BrickData @@ -85,6 +88,9 @@ where pub(crate) octree_size: u32, pub(crate) nodes: ObjectPool>, pub(crate) node_children: Vec>, // Children index values of each Node + + #[cfg(feature = "bevy_wgpu")] + pub(crate) views: Vec>, } #[derive(Default, Clone, Copy, Debug, PartialEq, Eq, Hash)]