From af1ae63dee3f77a298766e102f255daf8ea0004b Mon Sep 17 00:00:00 2001 From: davids91 Date: Thu, 11 Jul 2024 19:02:30 +0200 Subject: [PATCH 1/5] Updated bevy module to not copy octree data to GPU in every loop --- assets/shaders/viewport_render.wgsl | 16 +++++---- examples/bevy_wgpu_render.rs | 14 ++++++-- src/octree/raytracing/bevy/data.rs | 20 +++-------- src/octree/raytracing/bevy/mod.rs | 52 ++++++++++++++++++++++++----- src/octree/raytracing/bevy/types.rs | 21 ++++++++---- 5 files changed, 82 insertions(+), 41 deletions(-) diff --git a/assets/shaders/viewport_render.wgsl b/assets/shaders/viewport_render.wgsl index 0fcdcbf..9ae8726 100644 --- a/assets/shaders/viewport_render.wgsl +++ b/assets/shaders/viewport_render.wgsl @@ -336,7 +336,6 @@ fn traverse_brick( ) -> BrickHit{ var result: BrickHit; result.hit = false; - let pos = ( point_in_ray_at_distance(ray, *ray_current_distance) - bounds.min_position @@ -469,7 +468,10 @@ fn get_by_ray(ray_: Line) -> OctreeRayIntersection{ ); node_stack_i = 1; } + + var i = 0.; while(0 < node_stack_i && node_stack_i < max_depth) { + i += 1.; var current_bounds = node_stack[node_stack_i - 1].bounds; var current_node = nodes[node_stack[node_stack_i - 1].node]; //!NOTE: should be const, but then it can not be indexed dynamically var target_octant = node_stack[node_stack_i - 1].target_octant; @@ -630,22 +632,22 @@ struct Viewport { fov: f32, } -@group(0) @binding(1) +@group(0) @binding(0) var output_texture: texture_storage_2d; -@group(0) @binding(2) +@group(0) @binding(1) var viewport: Viewport; -@group(0) @binding(3) +@group(1) @binding(0) var octreeMetaData: OctreeMetaData; -@group(0) @binding(4) +@group(1) @binding(1) var nodes: array; -@group(0) @binding(5) +@group(1) @binding(2) var children_buffer: array; -@group(0) @binding(6) +@group(1) @binding(3) var voxels: array; @compute @workgroup_size(8, 8, 1) diff --git a/examples/bevy_wgpu_render.rs b/examples/bevy_wgpu_render.rs index c1df140..eec7d70 100644 --- a/examples/bevy_wgpu_render.rs +++ b/examples/bevy_wgpu_render.rs @@ -31,6 +31,8 @@ fn main() { #[cfg(feature = "bevy_wgpu")] fn setup(mut commands: Commands, images: ResMut>) { + use shocovox_rs::octree::raytracing::bevy::create_viewing_glass; + let origin = Vec3::new( ARRAY_DIMENSION as f32 * 2., ARRAY_DIMENSION as f32 / 2., @@ -74,7 +76,10 @@ fn setup(mut commands: Commands, images: ResMut>) { }; tree.insert( &V3c::new(x, y, z), - Albedo::from(r | (g << 8) | (b << 16) | 0xFF000000), + Albedo::default() + .with_red(r as u8) + .with_green(g as u8) + .with_blue(b as u8), ) .ok() .unwrap(); @@ -82,7 +87,8 @@ fn setup(mut commands: Commands, images: ResMut>) { } } } - let viewing_glass = tree.create_bevy_view( + let render_data = tree.create_bevy_view(); + let viewing_glass = create_viewing_glass( &Viewport { direction: (Vec3::new(0., 0., 0.) - origin).normalize(), origin, @@ -92,6 +98,7 @@ fn setup(mut commands: Commands, images: ResMut>) { DISPLAY_RESOLUTION, images, ); + commands.spawn(SpriteBundle { sprite: Sprite { custom_size: Some(Vec2::new(1024., 768.)), @@ -101,6 +108,7 @@ fn setup(mut commands: Commands, images: ResMut>) { ..default() }); commands.spawn(Camera2dBundle::default()); + commands.insert_resource(render_data); commands.insert_resource(viewing_glass); } @@ -116,7 +124,7 @@ fn rotate_camera( mut viewing_glass: ResMut, ) { let angle = { - let addition = ARRAY_DIMENSION as f32 / 10.; + let addition = ARRAY_DIMENSION as f32 / 100000.; let angle = angles_query.single().yaw + addition; if angle < 360. { angle diff --git a/src/octree/raytracing/bevy/data.rs b/src/octree/raytracing/bevy/data.rs index 8c7f445..81edb7a 100644 --- a/src/octree/raytracing/bevy/data.rs +++ b/src/octree/raytracing/bevy/data.rs @@ -1,17 +1,12 @@ use crate::object_pool::empty_marker; use crate::octree::{ - raytracing::{ - bevy::create_ouput_texture, - bevy::types::{OctreeMetaData, ShocoVoxViewingGlass, SizedNode, Viewport, Voxelement}, - }, + raytracing::bevy::types::{OctreeMetaData, ShocoVoxRenderData, SizedNode, Voxelement}, types::{NodeChildrenArray, NodeContent}, Octree, VoxelData, }; -use bevy::{ - asset::Assets, ecs::system::ResMut, math::Vec3, render::color::Color, render::texture::Image, -}; +use bevy::{math::Vec3, render::color::Color}; impl Octree where @@ -49,12 +44,7 @@ where meta } - pub fn create_bevy_view( - &self, - viewport: &Viewport, - resolution: [u32; 2], - images: ResMut>, - ) -> ShocoVoxViewingGlass { + pub fn create_bevy_view(&self) -> ShocoVoxRenderData { let meta = OctreeMetaData { octree_size: self.octree_size, voxel_brick_dim: DIM as u32, @@ -115,9 +105,7 @@ where nodes.push(sized_node); } - ShocoVoxViewingGlass { - output_texture: create_ouput_texture(resolution, images), - viewport: *viewport, + ShocoVoxRenderData { meta, nodes, children_buffer, diff --git a/src/octree/raytracing/bevy/mod.rs b/src/octree/raytracing/bevy/mod.rs index 4ee321c..d4d75db 100644 --- a/src/octree/raytracing/bevy/mod.rs +++ b/src/octree/raytracing/bevy/mod.rs @@ -6,7 +6,7 @@ pub use crate::octree::raytracing::bevy::types::{ }; use crate::octree::raytracing::bevy::types::{ - ShocoVoxLabel, ShocoVoxRenderNode, ShocoVoxRenderPipeline, + ShocoVoxLabel, ShocoVoxRenderData, ShocoVoxRenderNode, ShocoVoxRenderPipeline, }; use bevy::{ @@ -32,18 +32,33 @@ use bevy::{ use std::borrow::Cow; +pub fn create_viewing_glass( + viewport: &Viewport, + resolution: [u32; 2], + images: ResMut>, +) -> ShocoVoxViewingGlass { + ShocoVoxViewingGlass { + output_texture: create_ouput_texture(resolution, images), + viewport: *viewport, + } +} + impl FromWorld for ShocoVoxRenderPipeline { fn from_world(world: &mut World) -> Self { let render_device = world.resource::(); let viewing_glass_bind_group_layout = ShocoVoxViewingGlass::bind_group_layout(render_device); + let render_data_bind_group_layout = ShocoVoxRenderData::bind_group_layout(render_device); let shader = world .resource::() .load("shaders/viewport_render.wgsl"); let pipeline_cache = world.resource::(); let update_pipeline = pipeline_cache.queue_compute_pipeline(ComputePipelineDescriptor { label: None, - layout: vec![viewing_glass_bind_group_layout.clone()], + layout: vec![ + viewing_glass_bind_group_layout.clone(), + render_data_bind_group_layout.clone(), + ], push_constant_ranges: Vec::new(), shader, shader_defs: vec![], @@ -51,19 +66,23 @@ impl FromWorld for ShocoVoxRenderPipeline { }); ShocoVoxRenderPipeline { + update_tree: true, viewing_glass_bind_group_layout, + render_data_bind_group_layout, update_pipeline, - bind_group: None, + viewing_glass_bind_group: None, + tree_bind_group: None, } } } -fn prepare_bind_group( +fn prepare_bind_groups( gpu_images: Res>, fallback_image: Res, render_device: Res, mut pipeline: ResMut, octree_viewing_glass: Res, + render_data: Res, ) { let bind_group = octree_viewing_glass .as_bind_group( @@ -74,7 +93,21 @@ fn prepare_bind_group( ) .ok() .unwrap(); - pipeline.bind_group = Some(bind_group.bind_group); + pipeline.viewing_glass_bind_group = Some(bind_group.bind_group); + + if pipeline.update_tree { + let tree_bind_group = render_data + .as_bind_group( + &pipeline.render_data_bind_group_layout, + &render_device, + &gpu_images, + &fallback_image, + ) + .ok() + .unwrap(); + pipeline.tree_bind_group = Some(tree_bind_group.bind_group); + pipeline.update_tree = false; + } } pub(crate) fn create_ouput_texture( @@ -100,10 +133,11 @@ 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()); let render_app = app.sub_app_mut(RenderApp); render_app.add_systems( Render, - prepare_bind_group.in_set(RenderSet::PrepareBindGroups), + prepare_bind_groups.in_set(RenderSet::PrepareBindGroups), ); let mut render_graph = render_app.world.resource_mut::(); @@ -127,12 +161,13 @@ const WORKGROUP_SIZE: u32 = 8; impl render_graph::Node for ShocoVoxRenderNode { fn update(&mut self, world: &mut World) { let pipeline = world.resource::(); + let render_data = world.get_resource::(); let pipeline_cache = world.resource::(); if !self.ready { if let CachedPipelineState::Ok(_) = pipeline_cache.get_compute_pipeline_state(pipeline.update_pipeline) { - self.ready = true; + self.ready = render_data.is_some(); } } } @@ -151,7 +186,8 @@ impl render_graph::Node for ShocoVoxRenderNode { .command_encoder() .begin_compute_pass(&ComputePassDescriptor::default()); - pass.set_bind_group(0, pipeline.bind_group.as_ref().unwrap(), &[]); + pass.set_bind_group(0, pipeline.viewing_glass_bind_group.as_ref().unwrap(), &[]); + pass.set_bind_group(1, pipeline.tree_bind_group.as_ref().unwrap(), &[]); let pipeline = pipeline_cache .get_compute_pipeline(pipeline.update_pipeline) .unwrap(); diff --git a/src/octree/raytracing/bevy/types.rs b/src/octree/raytracing/bevy/types.rs index 4c6f3c0..cc9a23b 100644 --- a/src/octree/raytracing/bevy/types.rs +++ b/src/octree/raytracing/bevy/types.rs @@ -77,30 +77,37 @@ pub struct ShocoVoxRenderPlugin { #[derive(Resource, Clone, AsBindGroup, TypePath, ExtractResource)] #[type_path = "shocovox::gpu::ShocoVoxViewingGlass"] pub struct ShocoVoxViewingGlass { - #[storage_texture(1, image_format = Rgba8Unorm, access = ReadWrite)] + #[storage_texture(0, image_format = Rgba8Unorm, access = ReadWrite)] pub output_texture: Handle, - #[uniform(2, visibility(compute))] + #[uniform(1, visibility(compute))] pub viewport: Viewport, +} - #[uniform(3, visibility(compute))] +#[derive(Resource, Clone, AsBindGroup, TypePath, ExtractResource)] +#[type_path = "shocovox::gpu::ShocoVoxRenderData"] +pub struct ShocoVoxRenderData { + #[uniform(0, visibility(compute))] pub(crate) meta: OctreeMetaData, - #[storage(4, visibility(compute))] + #[storage(1, visibility(compute))] pub(crate) nodes: Vec, - #[storage(5, visibility(compute))] + #[storage(2, visibility(compute))] pub(crate) children_buffer: Vec, - #[storage(6, visibility(compute))] + #[storage(3, visibility(compute))] pub(crate) voxels: Vec, } #[derive(Resource)] pub(crate) struct ShocoVoxRenderPipeline { + pub update_tree: bool, pub(crate) viewing_glass_bind_group_layout: BindGroupLayout, + pub(crate) render_data_bind_group_layout: BindGroupLayout, pub(crate) update_pipeline: CachedComputePipelineId, - pub(crate) bind_group: Option, + pub(crate) viewing_glass_bind_group: Option, + pub(crate) tree_bind_group: Option, } #[derive(Debug, Hash, PartialEq, Eq, Clone, RenderLabel)] From d3f96d11a009161fa00291030817c3645be5ffaf Mon Sep 17 00:00:00 2001 From: davids91 Date: Fri, 12 Jul 2024 15:43:30 +0200 Subject: [PATCH 2/5] Reverted unneeded changes --- Cargo.toml | 8 +- ...wport_render.wgsl => viewport_render.wgsl} | 2 - examples/bevy_wgpu_render.rs | 17 +- examples/cpu_render.rs | 8 +- examples/wgpu_render.rs | 142 --- src/octree/detail.rs | 2 +- src/octree/raytracing/bevy/mod.rs | 2 +- src/octree/raytracing/mod.rs | 10 +- src/octree/raytracing/wgpu/data.rs | 259 ------ src/octree/raytracing/wgpu/mod.rs | 170 ---- src/octree/raytracing/wgpu/pipeline.rs | 246 ------ src/octree/raytracing/wgpu/raytracing.wgsl | 833 ------------------ src/octree/raytracing/wgpu/render.rs | 66 -- src/octree/raytracing/wgpu/types.rs | 110 --- src/octree/tests.rs | 1 + src/octree/types.rs | 74 +- src/spatial/math/tests.rs | 2 +- src/spatial/math/vector.rs | 8 +- 18 files changed, 59 insertions(+), 1901 deletions(-) rename assets/shaders/{bevy_viewport_render.wgsl => viewport_render.wgsl} (99%) delete mode 100644 examples/wgpu_render.rs delete mode 100644 src/octree/raytracing/wgpu/data.rs delete mode 100644 src/octree/raytracing/wgpu/mod.rs delete mode 100644 src/octree/raytracing/wgpu/pipeline.rs delete mode 100644 src/octree/raytracing/wgpu/raytracing.wgsl delete mode 100644 src/octree/raytracing/wgpu/render.rs delete mode 100644 src/octree/raytracing/wgpu/types.rs diff --git a/Cargo.toml b/Cargo.toml index 733f18b..9342fc5 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -9,19 +9,13 @@ license = "MIT OR Apache-2.0" default = [] raytracing = ["dep:image", "dep:show-image"] serialization = ["dep:serde"] -wgpu = ["raytracing", "dep:encase","dep:wgpu", "dep:winit", "dep:env_logger", "dep:futures"] -bevy_wgpu = ["raytracing", "dep:bevy"] +bevy_wgpu = ["raytracing", "dep:bevy", "dep:encase"] [dependencies] serde = { version = "1.0.183", features = ["derive"], optional = true } bendy = { git = "https://github.com/davids91/bendy.git" , features = ["std", "serde"]} array-init = "2.1.0" -wgpu = { version = "0.20.1", optional = true } -winit = { version = "0.30.3", optional = true } -env_logger = { version = "0.11.3", optional = true } -futures = { version = "0.3.30", optional = true } encase = { version = "0.9.0", optional = true } - # for example cpu_render image = { version = "0.25.1", optional = true } show-image = { version = "0.14.0", optional = true } diff --git a/assets/shaders/bevy_viewport_render.wgsl b/assets/shaders/viewport_render.wgsl similarity index 99% rename from assets/shaders/bevy_viewport_render.wgsl rename to assets/shaders/viewport_render.wgsl index 7d12385..9ae8726 100644 --- a/assets/shaders/bevy_viewport_render.wgsl +++ b/assets/shaders/viewport_render.wgsl @@ -688,8 +688,6 @@ fn update( textureStore(output_texture, pixel_location, vec4f(rgb_result, 1.)); } - - // Note: should be const var OCTANT_STEP_RESULT_LUT: array, 3>, 3> = array, 3>, 3>( array, 3>( diff --git a/examples/bevy_wgpu_render.rs b/examples/bevy_wgpu_render.rs index eec7d70..3c3de25 100644 --- a/examples/bevy_wgpu_render.rs +++ b/examples/bevy_wgpu_render.rs @@ -60,26 +60,23 @@ fn setup(mut commands: Commands, images: ResMut>) { && (ARRAY_DIMENSION / 2) <= z) { let r = if 0 == x % (ARRAY_DIMENSION / 4) { - (x as f32 / ARRAY_DIMENSION as f32 * 255.) as u32 + x as f32 / ARRAY_DIMENSION as f32 } else { - 128 + 0.5 }; let g = if 0 == y % (ARRAY_DIMENSION / 4) { - (y as f32 / ARRAY_DIMENSION as f32 * 255.) as u32 + y as f32 / ARRAY_DIMENSION as f32 } else { - 128 + 0.5 }; let b = if 0 == z % (ARRAY_DIMENSION / 4) { - (z as f32 / ARRAY_DIMENSION as f32 * 255.) as u32 + z as f32 / ARRAY_DIMENSION as f32 } else { - 128 + 0.5 }; tree.insert( &V3c::new(x, y, z), - Albedo::default() - .with_red(r as u8) - .with_green(g as u8) - .with_blue(b as u8), + Albedo::default().with_red(r).with_green(g).with_blue(b), ) .ok() .unwrap(); diff --git a/examples/cpu_render.rs b/examples/cpu_render.rs index e90ab48..482599f 100644 --- a/examples/cpu_render.rs +++ b/examples/cpu_render.rs @@ -30,10 +30,10 @@ fn main() { tree.insert( &V3c::new(x, y, z), Albedo::default() - .with_red((255 as f32 * x as f32 / TREE_SIZE as f32) as u8) - .with_green((255 as f32 * y as f32 / TREE_SIZE as f32) as u8) - .with_blue((255 as f32 * z as f32 / TREE_SIZE as f32) as u8) - .with_alpha(255), + .with_red(x as f32 / TREE_SIZE as f32) + .with_green(y as f32 / TREE_SIZE as f32) + .with_blue(z as f32 / TREE_SIZE as f32) + .with_alpha(1.), ) .ok() .unwrap(); diff --git a/examples/wgpu_render.rs b/examples/wgpu_render.rs deleted file mode 100644 index ae46de6..0000000 --- a/examples/wgpu_render.rs +++ /dev/null @@ -1,142 +0,0 @@ -use core::sync::atomic::Ordering; -use shocovox_rs::octree::raytracing::wgpu::SvxRenderApp; -use shocovox_rs::octree::Albedo; -use shocovox_rs::octree::V3c; -use shocovox_rs::octree::V3cf32; -use std::sync::Arc; -use winit::application::ApplicationHandler; -use winit::event::WindowEvent; -use winit::event_loop::ActiveEventLoop; -use winit::event_loop::{ControlFlow, EventLoop}; -use winit::window::Window; -use winit::window::WindowId; - -#[cfg(feature = "wgpu")] -const DISPLAY_RESOLUTION: [u32; 2] = [1024, 768]; - -#[cfg(feature = "wgpu")] -const ARRAY_DIMENSION: u32 = 128; - -#[cfg(feature = "wgpu")] -struct SvxRenderExample { - app: SvxRenderApp, - window: Option>, -} - -#[cfg(feature = "wgpu")] -impl ApplicationHandler for SvxRenderExample { - fn resumed(&mut self, event_loop: &ActiveEventLoop) { - self.window = Some(Arc::new( - event_loop - .create_window( - Window::default_attributes() - .with_title("Voxel Raytracing Render") - .with_inner_size(winit::dpi::PhysicalSize::new( - self.app.output_width(), - self.app.output_height(), - )), - ) - .unwrap(), - )); - futures::executor::block_on( - self.app - .rebuild_pipeline(self.window.as_ref().unwrap().clone()), - ) - } - - fn window_event(&mut self, event_loop: &ActiveEventLoop, _id: WindowId, event: WindowEvent) { - match event { - WindowEvent::CloseRequested => { - event_loop.exit(); - } - WindowEvent::RedrawRequested => { - self.app.execute_pipeline(); - self.window.as_ref().unwrap().request_redraw(); - } - WindowEvent::Resized(size) => { - futures::executor::block_on(self.app.set_output_size( - size.width, - size.height, - self.window.as_ref().unwrap().clone(), - )); - } - WindowEvent::KeyboardInput { - device_id: _, - event, - is_synthetic: _, - } => { - if let winit::keyboard::Key::Named(named) = event.logical_key { - if matches!(named, winit::keyboard::NamedKey::ArrowUp) { - self.app.update_viewport_origin(V3cf32::new(0.002, 0., 0.)); - } - } - } - _ => (), - } - } -} - -#[cfg(feature = "wgpu")] -fn main() { - // fill octree with data - let mut tree = shocovox_rs::octree::Octree::::new(ARRAY_DIMENSION) - .ok() - .unwrap(); - - tree.insert(&V3c::new(1, 3, 3), Albedo::from(0x66FFFF)) - .ok() - .unwrap(); - for x in 0..ARRAY_DIMENSION { - for y in 0..ARRAY_DIMENSION { - for z in 0..ARRAY_DIMENSION { - if ((x < (ARRAY_DIMENSION / 4) - || y < (ARRAY_DIMENSION / 4) - || z < (ARRAY_DIMENSION / 4)) - && (0 == x % 2 && 0 == y % 4 && 0 == z % 2)) - || ((ARRAY_DIMENSION / 2) <= x - && (ARRAY_DIMENSION / 2) <= y - && (ARRAY_DIMENSION / 2) <= z) - { - let r = if 0 == x % (ARRAY_DIMENSION / 4) { - (x as f32 / ARRAY_DIMENSION as f32 * 255.) as u32 - } else { - 128 - }; - let g = if 0 == y % (ARRAY_DIMENSION / 4) { - (y as f32 / ARRAY_DIMENSION as f32 * 255.) as u32 - } else { - 128 - }; - let b = if 0 == z % (ARRAY_DIMENSION / 4) { - (z as f32 / ARRAY_DIMENSION as f32 * 255.) as u32 - } else { - 128 - }; - tree.insert( - &V3c::new(x, y, z), - Albedo::from(r | (g << 8) | (b << 16) | 0xFF000000), - ) - .ok() - .unwrap(); - } - } - } - } - let showcase = Arc::new(tree); - - // Fire up the display - let event_loop = EventLoop::new().unwrap(); - event_loop.set_control_flow(ControlFlow::Poll); - let app = SvxRenderApp::new(DISPLAY_RESOLUTION[0], DISPLAY_RESOLUTION[1]); - let mut example = SvxRenderExample { app, window: None }; - //showcase.upload_to(&mut app); - - env_logger::init(); - let _ = event_loop.run_app(&mut example); -} - -#[cfg(not(feature = "wgpu"))] -fn main() { - println!("You probably forgot to enable the wgpu feature!"); - //nothing to do when the feature is not enabled -} diff --git a/src/octree/detail.rs b/src/octree/detail.rs index 754c9fe..daeda4e 100644 --- a/src/octree/detail.rs +++ b/src/octree/detail.rs @@ -94,7 +94,7 @@ where } } - #[cfg(any(feature = "bevy_wgpu", feature = "wgpu"))] + #[cfg(feature = "bevy_wgpu")] pub(in crate::octree) fn get_full(&self) -> [T; 8] { match &self.content { NodeChildrenArray::Children(c) => c.clone(), diff --git a/src/octree/raytracing/bevy/mod.rs b/src/octree/raytracing/bevy/mod.rs index d95dc89..d4d75db 100644 --- a/src/octree/raytracing/bevy/mod.rs +++ b/src/octree/raytracing/bevy/mod.rs @@ -51,7 +51,7 @@ impl FromWorld for ShocoVoxRenderPipeline { let render_data_bind_group_layout = ShocoVoxRenderData::bind_group_layout(render_device); let shader = world .resource::() - .load("shaders/bevy_viewport_render.wgsl"); + .load("shaders/viewport_render.wgsl"); let pipeline_cache = world.resource::(); let update_pipeline = pipeline_cache.queue_compute_pipeline(ComputePipelineDescriptor { label: None, diff --git a/src/octree/raytracing/mod.rs b/src/octree/raytracing/mod.rs index e9b4ac0..f54d278 100644 --- a/src/octree/raytracing/mod.rs +++ b/src/octree/raytracing/mod.rs @@ -1,16 +1,10 @@ pub mod raytracing_on_cpu; mod tests; -pub use crate::spatial::raytracing::Ray; - #[cfg(feature = "bevy_wgpu")] pub mod bevy; +pub use crate::spatial::raytracing::Ray; + #[cfg(feature = "bevy_wgpu")] pub use bevy::types::{ShocoVoxRenderPlugin, ShocoVoxViewingGlass, Viewport}; - -#[cfg(feature = "wgpu")] -pub mod wgpu; - -#[cfg(feature = "wgpu")] -pub use wgpu::types::{SvxRenderApp, Viewport}; diff --git a/src/octree/raytracing/wgpu/data.rs b/src/octree/raytracing/wgpu/data.rs deleted file mode 100644 index 9809fcc..0000000 --- a/src/octree/raytracing/wgpu/data.rs +++ /dev/null @@ -1,259 +0,0 @@ -use encase::StorageBuffer; -use std::num::NonZero; - -use crate::octree::{ - empty_marker, raytracing::wgpu::types::Voxelement, types::NodeChildrenArray, NodeContent, - Octree, VoxelData, -}; -use wgpu::util::DeviceExt; - -use super::{ - types::{OctreeMetaData, SizedNode}, - SvxRenderApp, -}; - -impl From<&Octree> for OctreeMetaData -where - T: Default + Clone + VoxelData, -{ - fn from(tree: &Octree) -> Self { - OctreeMetaData { - octree_size: tree.octree_size, - voxel_brick_dim: DIM as u32, - ambient_light_color: [1., 1., 1.].into(), - ambient_light_position: [DIM as f32, DIM as f32, DIM as f32].into(), - } - } -} - -impl Octree -where - T: Default + Clone + VoxelData, -{ - fn meta_set_is_leaf(sized_node_meta: &mut u32, is_leaf: bool) { - *sized_node_meta = - (*sized_node_meta & 0x00FFFFFF) | if is_leaf { 0x01000000 } else { 0x00000000 }; - } - - fn meta_set_node_occupancy_bitmap(sized_node_meta: &mut u32, bitmap: u8) { - *sized_node_meta = (*sized_node_meta & 0xFFFFFF00) | bitmap as u32; - } - - fn create_meta(&self, node_key: usize) -> u32 { - let node = self.nodes.get(node_key); - let mut meta = 0; - match node { - NodeContent::Leaf(_) => { - Self::meta_set_is_leaf(&mut meta, true); - Self::meta_set_node_occupancy_bitmap( - &mut meta, - self.occupied_8bit(node_key as u32), - ); - } - NodeContent::Internal(occupied_bits) => { - Self::meta_set_is_leaf(&mut meta, false); - Self::meta_set_node_occupancy_bitmap(&mut meta, *occupied_bits); - } - _ => { - Self::meta_set_is_leaf(&mut meta, false); - Self::meta_set_node_occupancy_bitmap(&mut meta, 0x00); - } - }; - meta - } - - pub(crate) fn upload_to(&self, app: &mut SvxRenderApp) { - // parse octree - let mut nodes = Vec::new(); - let mut children = Vec::new(); - let mut voxels = Vec::new(); - - for i in 0..self.nodes.len() { - if !self.nodes.key_is_valid(i) { - continue; - } - let mut sized_node = SizedNode { - sized_node_meta: self.create_meta(i), - children_start_at: children.len() as u32, - voxels_start_at: empty_marker(), - }; - if let NodeContent::Leaf(data) = self.nodes.get(i) { - debug_assert!(matches!( - self.node_children[i].content, - NodeChildrenArray::OccupancyBitmap(_) - )); - let occupied_bits = match self.node_children[i].content { - NodeChildrenArray::OccupancyBitmap(bitmap) => bitmap, - _ => panic!("Found Leaf Node without occupancy bitmap!"), - }; - children.extend_from_slice(&[ - (occupied_bits & 0x00000000FFFFFFFF) as u32, - ((occupied_bits & 0xFFFFFFFF00000000) >> 32) as u32, - ]); - sized_node.voxels_start_at = voxels.len() as u32; - for z in 0..DIM { - for y in 0..DIM { - for x in 0..DIM { - let albedo = data[x][y][z].albedo(); - let content = data[x][y][z].user_data(); - voxels.push(Voxelement { albedo, content }) - } - } - } - } else { - //Internal nodes - children.extend_from_slice(&self.node_children[i].get_full()); - } - nodes.push(sized_node); - } - - debug_assert!(0 < nodes.len()); - debug_assert!(0 < children.len()); - debug_assert!(0 < voxels.len()); - - // Create bind group layout - let layout = app - .device - .as_ref() - .expect("Expected Render Device") - .create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { - entries: &[ - // metadata - wgpu::BindGroupLayoutEntry { - binding: 0, - visibility: wgpu::ShaderStages::COMPUTE, - ty: wgpu::BindingType::Buffer { - ty: wgpu::BufferBindingType::Uniform, - has_dynamic_offset: false, - min_binding_size: None, - }, - count: None, - }, - // nodes - wgpu::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: Some(NonZero::new(nodes.len() as u32).unwrap()), - }, - // children - wgpu::BindGroupLayoutEntry { - binding: 2, - visibility: wgpu::ShaderStages::COMPUTE, - ty: wgpu::BindingType::Buffer { - ty: wgpu::BufferBindingType::Storage { read_only: true }, - has_dynamic_offset: false, - min_binding_size: None, - }, - count: Some(NonZero::new(children.len() as u32).unwrap()), - }, - // voxels - wgpu::BindGroupLayoutEntry { - binding: 3, - visibility: wgpu::ShaderStages::COMPUTE, - ty: wgpu::BindingType::Buffer { - ty: wgpu::BufferBindingType::Storage { read_only: true }, - has_dynamic_offset: false, - min_binding_size: None, - }, - count: Some(NonZero::new(voxels.len() as u32).unwrap()), - }, - ], - label: Some("Octree_Layout"), - }); - - // Upload data to buffers - let octree_meta = OctreeMetaData::from(self); - let mut buffer = StorageBuffer::new(Vec::::new()); - buffer.write(&octree_meta).unwrap(); - let metadata_buffer = app - .device - .as_ref() - .expect("Expected SvxRenderApp to have a vaild device!") - .create_buffer_init(&wgpu::util::BufferInitDescriptor { - label: Some("Octree Metadata Buffer"), - contents: &buffer.into_inner(), - usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, - }); - let mut buffer = StorageBuffer::new(Vec::::new()); - buffer.write(&nodes).unwrap(); - let nodes_buffer = app - .device - .as_ref() - .expect("Expected SvxRenderApp to have a vaild device!") - .create_buffer_init(&wgpu::util::BufferInitDescriptor { - label: Some("Octree Metadata Buffer"), - contents: &buffer.into_inner(), - usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, - }); - - let mut buffer = StorageBuffer::new(Vec::::new()); - buffer.write(&children).unwrap(); - let children_buffer = app - .device - .as_ref() - .expect("Expected SvxRenderApp to have a vaild device!") - .create_buffer_init(&wgpu::util::BufferInitDescriptor { - label: Some("Octree Metadata Buffer"), - contents: &buffer.into_inner(), - usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, - }); - - let mut buffer = StorageBuffer::new(Vec::::new()); - buffer.write(&voxels).unwrap(); - let voxels_buffer = app - .device - .as_ref() - .expect("Expected SvxRenderApp to have a vaild device!") - .create_buffer_init(&wgpu::util::BufferInitDescriptor { - label: Some("Octree Metadata Buffer"), - contents: &buffer.into_inner(), - usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, - }); - - // Create bind group - let group = app - .device - .as_ref() - .expect("Expected SvxRenderApp to have a vaild device!") - .create_bind_group(&wgpu::BindGroupDescriptor { - layout: &layout, - entries: &[ - // wgpu::BindGroupEntry { - // binding: 0, - // resource: ??.as_entire_binding(), - // }, - wgpu::BindGroupEntry { - binding: 0, - resource: metadata_buffer.as_entire_binding(), - }, - wgpu::BindGroupEntry { - binding: 1, - resource: nodes_buffer.as_entire_binding(), - }, - wgpu::BindGroupEntry { - binding: 2, - resource: children_buffer.as_entire_binding(), - }, - wgpu::BindGroupEntry { - binding: 3, - resource: voxels_buffer.as_entire_binding(), - }, - ], - label: Some("camera_bind_group"), - }); - let render_pipeline_layout = app - .device - .as_ref() - .expect("Expected SvxRenderApp to have a vaild device!") - .create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { - label: Some("Render Pipeline Layout"), - bind_group_layouts: &[&layout], - push_constant_ranges: &[], - }); - } -} diff --git a/src/octree/raytracing/wgpu/mod.rs b/src/octree/raytracing/wgpu/mod.rs deleted file mode 100644 index 25023e6..0000000 --- a/src/octree/raytracing/wgpu/mod.rs +++ /dev/null @@ -1,170 +0,0 @@ -mod data; -mod pipeline; -mod render; -pub mod types; - -pub use crate::octree::raytracing::{SvxRenderApp, Viewport}; -use crate::spatial::math::vector::V3cf32; -use encase::UniformBuffer; -use std::sync::Arc; -use winit::window::Window; - -impl<'a> SvxRenderApp { - pub fn new(output_width: u32, output_height: u32) -> Self { - Self { - viewport: Viewport::default(), - output_width, - output_height, - texture_extent: wgpu::Extent3d { - width: output_width, - height: output_height, - depth_or_array_layers: 1, - }, - can_render: false.into(), - wgpu_instance: wgpu::Instance::default(), - adapter: None, - surface: None, - device: None, - pipeline: None, - queue: None, - dynamic_group: None, - output_texture: None, - output_texture_render: None, - viewport_buffer: None, - } - } - - pub fn output_width(&self) -> u32 { - self.output_width - } - - pub fn output_height(&self) -> u32 { - self.output_height - } - - pub async fn set_output_size(&mut self, width: u32, height: u32, window: Arc) { - self.output_width = width; - self.output_height = height; - self.rebuild_pipeline(window).await; - } - - pub fn viewport(&self) -> &Viewport { - &self.viewport - } - - pub fn set_viewport(&mut self, viewport: Viewport) { - if viewport == self.viewport { - return; - } - - let mut buffer = UniformBuffer::new(Vec::::new()); - buffer.write(&viewport).unwrap(); - self.viewport = viewport; - self.queue - .as_ref() - .expect("Expected SvxRenderApp to have a vaild rendering queue!") - .write_buffer( - self.viewport_buffer - .as_ref() - .expect("Expected SvxRenderApp to have a vaild Viewport buffer!"), - 0, - &buffer.into_inner(), - ) - } - - pub fn set_viewport_origin(&mut self, origin: V3cf32) { - self.viewport.origin = origin; - let mut buffer = UniformBuffer::new(Vec::::new()); - buffer.write(&self.viewport).unwrap(); - self.queue - .as_ref() - .expect("Expected SvxRenderApp to have a vaild rendering queue!") - .write_buffer( - self.viewport_buffer - .as_ref() - .expect("Expected SvxRenderApp to have a vaild Viewport buffer!"), - 0, - &buffer.into_inner(), - ) - } - - pub fn update_viewport_origin(&mut self, delta: V3cf32) { - self.viewport.origin = self.viewport.origin + delta; - let mut buffer = UniformBuffer::new(Vec::::new()); - buffer.write(&self.viewport).unwrap(); - self.queue - .as_ref() - .expect("Expected SvxRenderApp to have a vaild rendering queue!") - .write_buffer( - self.viewport_buffer - .as_ref() - .expect("Expected SvxRenderApp to have a vaild Viewport buffer!"), - 0, - &buffer.into_inner(), - ) - } - - pub fn update_viewport_direction(&mut self, direction: V3cf32) { - self.viewport.direction = direction; - let mut buffer = UniformBuffer::new(Vec::::new()); - buffer.write(&self.viewport).unwrap(); - self.queue - .as_ref() - .expect("Expected SvxRenderApp to have a vaild rendering queue!") - .write_buffer( - self.viewport_buffer - .as_ref() - .expect("Expected SvxRenderApp to have a vaild Viewport buffer!"), - 0, - &buffer.into_inner(), - ) - } - - pub fn update_viewport_glass_width(&mut self, width: f32) { - self.viewport.w_h_fov.x = width; - let mut buffer = UniformBuffer::new(Vec::::new()); - buffer.write(&self.viewport).unwrap(); - self.queue - .as_ref() - .expect("Expected SvxRenderApp to have a vaild rendering queue!") - .write_buffer( - self.viewport_buffer - .as_ref() - .expect("Expected SvxRenderApp to have a vaild Viewport buffer!"), - 0, - &buffer.into_inner(), - ) - } - - pub fn update_viewport_glass_height(&mut self, height: f32) { - self.viewport.w_h_fov.y = height; - let mut buffer = UniformBuffer::new(Vec::::new()); - buffer.write(&self.viewport).unwrap(); - self.queue - .as_ref() - .expect("Expected SvxRenderApp to have a vaild rendering queue!") - .write_buffer( - self.viewport_buffer - .as_ref() - .expect("Expected SvxRenderApp to have a vaild Viewport buffer!"), - 0, - &buffer.into_inner(), - ) - } - - pub fn update_viewport_glass_fov(&mut self, fov: f32) { - self.viewport.w_h_fov.z = fov; - let mut buffer = UniformBuffer::new(Vec::::new()); - buffer.write(&self.viewport).unwrap(); - self.queue - .as_ref() - .expect("Expected SvxRenderApp to have a vaild rendering queue!") - .write_buffer( - self.viewport_buffer - .as_ref() - .expect("Expected SvxRenderApp to have a vaild Viewport buffer!"), - 0, - &buffer.into_inner(), - ) - } -} diff --git a/src/octree/raytracing/wgpu/pipeline.rs b/src/octree/raytracing/wgpu/pipeline.rs deleted file mode 100644 index 1c85a54..0000000 --- a/src/octree/raytracing/wgpu/pipeline.rs +++ /dev/null @@ -1,246 +0,0 @@ -use crate::octree::raytracing::wgpu::SvxRenderApp; -use encase::UniformBuffer; -use std::borrow::Cow; -use std::sync::Arc; -use wgpu::{util::DeviceExt, TextureUsages}; -use winit::window::Window; - -impl SvxRenderApp { - pub async fn rebuild_pipeline(&mut self, window: Arc) { - let surface = self.wgpu_instance.create_surface(window).unwrap(); - - let adapter = self - .wgpu_instance - .request_adapter(&wgpu::RequestAdapterOptions { - power_preference: wgpu::PowerPreference::default(), - force_fallback_adapter: false, - // Request an adapter which can render to our surface - compatible_surface: Some(&surface), - }) - .await - .expect("Failed to find an appropriate adapter"); - - let (device, queue) = adapter - .request_device( - &wgpu::DeviceDescriptor { - label: None, - required_features: wgpu::Features::TEXTURE_ADAPTER_SPECIFIC_FORMAT_FEATURES, - // Make sure we use the texture resolution limits from the adapter, so we can support images the size of the swapchain. - required_limits: wgpu::Limits::downlevel_defaults() - .using_resolution(adapter.limits()), - }, - None, - ) - .await - .expect("Failed to create device"); - - // Create output texture and fill it with a single color - let output_texture = device.create_texture(&wgpu::TextureDescriptor { - size: self.texture_extent, - mip_level_count: 1, - sample_count: 1, - dimension: wgpu::TextureDimension::D2, - format: wgpu::TextureFormat::Rgba8Unorm, - usage: TextureUsages::STORAGE_BINDING - | TextureUsages::COPY_SRC - | TextureUsages::COPY_DST, - label: Some("output_texture"), - view_formats: &[], - }); - let output_texture_render = device.create_texture(&wgpu::TextureDescriptor { - size: self.texture_extent, - mip_level_count: 1, - sample_count: 1, - dimension: wgpu::TextureDimension::D2, - format: wgpu::TextureFormat::Rgba8Unorm, - usage: TextureUsages::TEXTURE_BINDING | TextureUsages::COPY_DST, - label: Some("output_texture"), - view_formats: &[], - }); - let mut default_image = image::DynamicImage::ImageRgba8(image::RgbaImage::new( - self.output_width, - self.output_height, - )); - - for x in 0..self.output_width { - for y in 0..self.output_height { - use image::GenericImage; - if 0 == x % 50 || 0 == y % 50 { - default_image.put_pixel(x, y, image::Rgba([200, 255, 200, 255])); - } else { - default_image.put_pixel( - x, - y, - image::Rgba([ - (255 * x / self.output_width) as u8, - 255, - (255 * y / self.output_width) as u8, - 255, - ]), - ); - } - } - } - - queue.write_texture( - wgpu::ImageCopyTexture { - texture: &output_texture, - mip_level: 0, - origin: wgpu::Origin3d::ZERO, - aspect: wgpu::TextureAspect::All, - }, - &default_image.to_rgba8(), - wgpu::ImageDataLayout { - offset: 0, - bytes_per_row: Some(4 * self.output_width), - rows_per_image: Some(self.output_height), - }, - self.texture_extent, - ); - - // create texture view and sampler - let output_texture_view = - output_texture.create_view(&wgpu::TextureViewDescriptor::default()); - let output_texture_render_view = - output_texture_render.create_view(&wgpu::TextureViewDescriptor::default()); - let output_texture_sampler = device.create_sampler(&wgpu::SamplerDescriptor { - address_mode_u: wgpu::AddressMode::ClampToEdge, - address_mode_v: wgpu::AddressMode::ClampToEdge, - address_mode_w: wgpu::AddressMode::ClampToEdge, - mag_filter: wgpu::FilterMode::Linear, - min_filter: wgpu::FilterMode::Nearest, - mipmap_filter: wgpu::FilterMode::Nearest, - ..Default::default() - }); - - // re-create viewport - let mut buffer = UniformBuffer::new(Vec::::new()); - buffer.write(&self.viewport).unwrap(); - let viewport_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor { - label: Some("Octree Metadata Buffer"), - contents: &buffer.into_inner(), - usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, - }); - - // Create bind group layout - let layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { - entries: &[ - // output_texture - wgpu::BindGroupLayoutEntry { - binding: 0, - visibility: wgpu::ShaderStages::COMPUTE | wgpu::ShaderStages::FRAGMENT, - ty: wgpu::BindingType::StorageTexture { - access: wgpu::StorageTextureAccess::ReadWrite, - format: wgpu::TextureFormat::Rgba8Unorm, - view_dimension: wgpu::TextureViewDimension::D2, - }, - count: None, - }, - // output_texture_render - wgpu::BindGroupLayoutEntry { - binding: 1, - visibility: wgpu::ShaderStages::COMPUTE | wgpu::ShaderStages::FRAGMENT, - ty: wgpu::BindingType::Texture { - multisampled: false, - view_dimension: wgpu::TextureViewDimension::D2, - sample_type: wgpu::TextureSampleType::Float { filterable: true }, - }, - count: None, - }, - // output_texture sampler - wgpu::BindGroupLayoutEntry { - binding: 2, - visibility: wgpu::ShaderStages::FRAGMENT, - ty: wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Filtering), - count: None, - }, - // viewport - wgpu::BindGroupLayoutEntry { - binding: 3, - visibility: wgpu::ShaderStages::COMPUTE | wgpu::ShaderStages::FRAGMENT, - ty: wgpu::BindingType::Buffer { - ty: wgpu::BufferBindingType::Uniform, - has_dynamic_offset: false, - min_binding_size: None, - }, - count: None, - }, - ], - label: Some("Dynamic_layout"), - }); - - let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor { - layout: &layout, - entries: &[ - wgpu::BindGroupEntry { - binding: 0, - resource: wgpu::BindingResource::TextureView(&output_texture_view), - }, - wgpu::BindGroupEntry { - binding: 1, - resource: wgpu::BindingResource::TextureView(&output_texture_render_view), - }, - wgpu::BindGroupEntry { - binding: 2, - resource: wgpu::BindingResource::Sampler(&output_texture_sampler), - }, - wgpu::BindGroupEntry { - binding: 3, - resource: viewport_buffer.as_entire_binding(), - }, - ], - label: Some("Dynamic_bind_group"), - }); - - let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { - label: None, - bind_group_layouts: &[&layout], - push_constant_ranges: &[], - }); - - let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor { - label: None, - source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(include_str!("raytracing.wgsl"))), - }); - - let swapchain_capabilities = surface.get_capabilities(&adapter); - let swapchain_format = swapchain_capabilities.formats[0]; - self.pipeline = Some( - device.create_render_pipeline(&wgpu::RenderPipelineDescriptor { - label: None, - layout: Some(&pipeline_layout), - vertex: wgpu::VertexState { - module: &shader, - entry_point: "vs_main", - buffers: &[], - compilation_options: Default::default(), - }, - fragment: Some(wgpu::FragmentState { - module: &shader, - entry_point: "fs_main", - compilation_options: Default::default(), - targets: &[Some(swapchain_format.into())], - }), - primitive: wgpu::PrimitiveState::default(), - depth_stencil: None, - multisample: wgpu::MultisampleState::default(), - multiview: None, - }), - ); - - let config = surface - .get_default_config(&adapter, self.output_width, self.output_height) - .unwrap(); - surface.configure(&device, &config); - - // Set object state - self.adapter = Some(adapter); - self.surface = Some(surface); - self.queue = Some(queue); - self.device = Some(device); - self.dynamic_group = Some(bind_group); - self.output_texture = Some(output_texture); - self.output_texture_render = Some(output_texture_render); - self.viewport_buffer = Some(viewport_buffer); - } -} diff --git a/src/octree/raytracing/wgpu/raytracing.wgsl b/src/octree/raytracing/wgpu/raytracing.wgsl deleted file mode 100644 index 38c23aa..0000000 --- a/src/octree/raytracing/wgpu/raytracing.wgsl +++ /dev/null @@ -1,833 +0,0 @@ -struct Line { - origin: vec3f, - direction: vec3f, -} - -struct Plane { - point: vec3f, - normal: vec3f, -} - -struct Cube { - min_position: vec3f, - size: f32, -} - -const FLOAT_ERROR_TOLERANCE = 0.00001; -const OOB_OCTANT = 8u; - -//crate::spatial::raytracing::Cube::contains_point -fn cube_contains_point(cube: Cube, p: vec3f) -> bool{ - let min_cn = p >= cube.min_position - FLOAT_ERROR_TOLERANCE; - let max_cn = p < (cube.min_position + cube.size + FLOAT_ERROR_TOLERANCE); - return ( - min_cn.x && min_cn.y && min_cn.z && max_cn.x && max_cn.y && max_cn.z - ); -} - -//Rust::unwrap_or -fn impact_or(impact: CubeRayIntersection, or: f32) -> f32{ - if(impact.hit && impact.impact_hit){ - return impact.impact_distance; - } - return or; -} - -//crate::spatial::math::hash_region -fn hash_region(offset: vec3f, size: f32) -> u32 { - let midpoint = vec3f(size / 2., size / 2., size / 2.); - return u32(offset.x >= midpoint.x) - + u32(offset.z >= midpoint.z) * 2u - + u32(offset.y >= midpoint.y) * 4u; -} - -//crate::spatial::math::offset_region -fn offset_region(octant: u32) -> vec3f { - switch(octant){ - case 0u { return vec3f(0., 0., 0.); } - case 1u { return vec3f(1., 0., 0.); } - case 2u { return vec3f(0., 0., 1.); } - case 3u { return vec3f(1., 0., 1.); } - case 4u { return vec3f(0., 1., 0.); } - case 5u { return vec3f(1., 1., 0.); } - case 6u { return vec3f(0., 1., 1.); } - case 7u, default { return vec3f(1.,1.,1.); } - } -} - -//crate::spatial::mod::Cube::child_bounds_for -fn child_bounds_for(bounds: Cube, octant: u32) -> Cube{ - var result: Cube; - let child_size = bounds.size / 2.; - result.min_position = bounds.min_position + (offset_region(octant) * child_size); - result.size = child_size; - return result; -} - -struct PlaneLineIntersection { - hit: bool, - d: f32, -} - -//crate::spatial::math::plane_line_intersection -fn plane_line_intersection(plane: Plane, line: Line) -> PlaneLineIntersection { - var result: PlaneLineIntersection; - let origins_diff = plane.point - line.origin; - let plane_line_dot_to_plane = dot(origins_diff, plane.normal); - let directions_dot = dot(line.direction, plane.normal); - - if 0. == directions_dot { - // line and plane is paralell - if 0. == dot(origins_diff, plane.normal) { - // The distance is zero because the origin is already on the plane - result.hit = true; - result.d = 0.; - } else { - result.hit = false; - } - } else { - result.hit = true; - result.d = plane_line_dot_to_plane / directions_dot; - } - return result; -} - -//crate::spatial::raytracing::Cube::face -fn get_cube_face(cube: Cube, face_index: u32) -> Plane{ - var result: Plane; - switch(face_index){ - case 0u { result.normal = vec3f(0.,0.,-1.); } - case 1u { result.normal = vec3f(-1.,0.,0.); } - case 2u { result.normal = vec3f(0.,0.,1.); } - case 3u { result.normal = vec3f(1.,0.,0.); } - case 4u { result.normal = vec3f(0.,1.,0.); } - case 5u, default { result.normal = vec3f(0.,-1.,0.); } - } - let half_size = cube.size / 2.; - let midpoint = cube.min_position + half_size; - result.point = midpoint + result.normal * half_size; - return result; -} - -struct CubeRayIntersection { - hit: bool, - impact_hit: bool, - impact_distance: f32, - exit_distance: f32, -} - -//crate::spatial::raytracing::Ray::point_at -fn point_in_ray_at_distance(ray: Line, d: f32) -> vec3f{ - return ray.origin + ray.direction * d; -} - -//crate::spatial::raytracing::Cube::intersect_ray -fn cube_intersect_ray(cube: Cube, ray: Line) -> CubeRayIntersection{ - var result: CubeRayIntersection; - let max_position = cube.min_position + vec3f(cube.size, cube.size, cube.size); - let t1 = (cube.min_position.x - ray.origin.x) / ray.direction.x; - let t2 = (max_position.x - ray.origin.x) / ray.direction.x; - let t3 = (cube.min_position.y - ray.origin.y) / ray.direction.y; - let t4 = (max_position.y - ray.origin.y) / ray.direction.y; - let t5 = (cube.min_position.z - ray.origin.z) / ray.direction.z; - let t6 = (max_position.z - ray.origin.z) / ray.direction.z; - - let tmin = max(max(min(t1, t2), min(t3, t4)), min(t5, t6)); - let tmax = min(min(max(t1, t2), max(t3, t4)), max(t5, t6)); - - if tmax < 0. || tmin > tmax{ - result.hit = false; - return result; - } - - let p = point_in_ray_at_distance(ray, tmin); - - - if tmin < 0.0 { - result.hit = true; - result.impact_hit = false; - result.exit_distance = tmax; - return result; - } - - result.hit = true; - result.impact_hit = true; - result.impact_distance = tmin; - result.exit_distance = tmax; - return result; -} - -fn cube_impact_normal(cube: Cube, impact_point: vec3f) -> vec3f{ - var impact_normal = vec3f(0.,0.,0.); - let mid_to_impact = cube.min_position + vec3f(cube.size / 2.) - impact_point; - let mid_to_impact_abs = abs(mid_to_impact); - let max_component = max( - mid_to_impact_abs.x, - max(mid_to_impact_abs.y, mid_to_impact_abs.z) - ); - if max_component - mid_to_impact_abs.x < FLOAT_ERROR_TOLERANCE { - impact_normal.x = -mid_to_impact.x; - } - if max_component - mid_to_impact_abs.y < FLOAT_ERROR_TOLERANCE { - impact_normal.y = -mid_to_impact.y; - } - if max_component - mid_to_impact_abs.z < FLOAT_ERROR_TOLERANCE { - impact_normal.z = -mid_to_impact.z; - } - return normalize(impact_normal); -} - -struct NodeStackItem { - bounds: Cube, - node: u32, - sized_node_meta: u32, - target_octant: u32, -} - -//crate::octree:raytracing::NodeStackItem::new -fn new_node_stack_item( - bounds: Cube, - node: u32, - sized_node_meta: u32, - target_octant: u32 -) -> NodeStackItem { - var result: NodeStackItem; - result.bounds = bounds; - result.node = node; - result.target_octant = target_octant; - result.sized_node_meta = sized_node_meta; - return result; -} - -//crate::octree:raytracing::get_dda_scale_factors -fn get_dda_scale_factors(ray: Line) -> vec3f { - return vec3f( - sqrt( - 1. - + pow(ray.direction.z / ray.direction.x, 2.) - + pow(ray.direction.y / ray.direction.x, 2.) - ), - sqrt( - pow(ray.direction.x / ray.direction.y, 2.) - + 1. - + pow(ray.direction.z / ray.direction.y, 2.) - ), - sqrt( - pow(ray.direction.x / ray.direction.z, 2.) - + pow(ray.direction.y / ray.direction.z, 2.) - + 1. - ), - ); -} - -//crate::octree::raytracing::dda_step_to_next_sibling -fn dda_step_to_next_sibling( - ray: Line, - ray_current_distance: ptr, - current_bounds: Cube, - ray_scale_factors: vec3f -) -> vec3f { - var signum_vec = sign(ray.direction); - let p = point_in_ray_at_distance(ray, *ray_current_distance); - let steps_needed = ( - p - current_bounds.min_position - - (current_bounds.size * max(sign(ray.direction), vec3f(0.,0.,0.))) - ); - - let d = ( - vec3f(*ray_current_distance, *ray_current_distance, *ray_current_distance) - + abs(steps_needed * ray_scale_factors) - ); - *ray_current_distance = min(d.x, min(d.y, d.z)); - - var result = vec3f(0., 0., 0.); - if abs(*ray_current_distance - d.x) < FLOAT_ERROR_TOLERANCE { - result.x = f32(abs(current_bounds.size)) * signum_vec.x; - } - if abs(*ray_current_distance - d.y) < FLOAT_ERROR_TOLERANCE { - result.y = f32(abs(current_bounds.size)) * signum_vec.y; - } - if abs(*ray_current_distance - d.z) < FLOAT_ERROR_TOLERANCE { - result.z = f32(abs(current_bounds.size)) * signum_vec.z; - } - return result; -} - -// Unique to this implementation, not adapted from rust code, corresponds to: -//crate::octree::raytracing::classic_raytracing_on_bevy_wgpu::meta_set_is_leaf -fn is_leaf(sized_node_meta: u32) -> bool { - return 0 < (0x01000000 & sized_node_meta); -} - -// Unique to this implementation, not adapted from rust code, corresponds to: -//crate::octree::raytracing::classic_raytracing_on_bevy_wgpu::meta_set_node_occupancy_bitmap -fn get_node_occupancy_bitmap(sized_node_meta: u32) -> u32 { - return (0x000000FF & sized_node_meta); -} - -//crate::spatial::math::step_octant -fn step_octant(octant: u32, step: vec3f) -> u32 { - let octant_pos_in_32bits = 4 * octant; - return ((OCTANT_STEP_RESULT_LUT[u32(sign(step.x) + 1)][u32(sign(step.y) + 1)][u32(sign(step.z) + 1)] - & (0x0Fu << octant_pos_in_32bits)) - >> octant_pos_in_32bits) & 0x0Fu; -} - -//crate::spatial::math::hash_direction -fn hash_direction(direction: vec3f) -> u32 { - let offset = vec3f(1.) + normalize(direction); - return hash_region(offset, 2.); -} - -// Functionality-wise this function is more generic, than its counterpart -// and is used in voxel brick mapping too -//crate::spatial::math::flat_projection -fn flat_projection(i: vec3u, dimensions: vec2u) -> u32 { - return (i.x + (i.y * dimensions.y) + (i.z * dimensions.x * dimensions.y)); -} - -//crate::spatial::math::position_in_bitmap_64bits -fn position_in_bitmap_64bits(i: vec3u, dimension: u32) -> u32{ - let pos_inside_bitmap_space = i * 4 / dimension; - //let pos_inside_bitmap_space = vec3u((vec3f(i) * 4.) / f32(dimension)); - let pos_inside_bitmap = flat_projection( - pos_inside_bitmap_space, vec2u(4, 4) - ); - return pos_inside_bitmap; -} - -// Unique to this implementation, not adapted from rust code -fn get_occupancy_in_bitmap_64bits( - bit_position: u32, - bitmap_lsb: u32, - bitmap_msb: u32 -) -> bool { - // not possible to create a position mask directly, because of missing u64 type - if bit_position < 32 { - let pos_mask = u32(0x01u << bit_position); - return 0 < (bitmap_lsb & pos_mask); - } - let pos_mask = u32(0x01u << (bit_position - 32)); - return 0 < (bitmap_msb & pos_mask); -} - -struct BrickHit{ - hit: bool, - index: vec3u -} - -fn traverse_brick( - ray: Line, - ray_current_distance: ptr, - brick_index_start: u32, - occupancy_bitmap_lsb: u32, - occupancy_bitmap_msb: u32, - bounds: Cube, - ray_scale_factors: vec3f, - direction_lut_index: u32, - unit_in_bitmap_space: f32, - dimension: u32 -) -> BrickHit{ - var result: BrickHit; - result.hit = false; - - let pos = ( - point_in_ray_at_distance(ray, *ray_current_distance) - - bounds.min_position - ); - var current_index = vec3i( - clamp(i32(pos.x), 0, i32(dimension - 1)), - clamp(i32(pos.y), 0, i32(dimension - 1)), - clamp(i32(pos.z), 0, i32(dimension - 1)) - ); - let brick_unit = bounds.size / f32(dimension); - var current_bounds = Cube( - bounds.min_position + vec3f(current_index) * brick_unit, - brick_unit - ); - - let start_pos_in_bitmap = position_in_bitmap_64bits(vec3u(current_index), dimension); - if ( - 0 == (RAY_TO_LEAF_OCCUPANCY_BITMASK_LUT[start_pos_in_bitmap][direction_lut_index * 2] - & occupancy_bitmap_lsb) - && 0 == (RAY_TO_LEAF_OCCUPANCY_BITMASK_LUT[start_pos_in_bitmap][direction_lut_index * 2 + 1] - & occupancy_bitmap_msb) - ){ - result.hit = false; - return result; - } - - var prev_bitmap_position_full_resolution = vec3u(vec3f(current_index) * unit_in_bitmap_space); - loop{ - if current_index.x < 0 - || current_index.x >= i32(dimension) - || current_index.y < 0 - || current_index.y >= i32(dimension) - || current_index.z < 0 - || current_index.z >= i32(dimension) - { - result.hit = false; - return result; - } - - let bitmap_position_full_resolution = vec3u(vec3f(current_index) * unit_in_bitmap_space); - let differs = bitmap_position_full_resolution != prev_bitmap_position_full_resolution; - if(differs.x || differs.y || differs.z) { - prev_bitmap_position_full_resolution = bitmap_position_full_resolution; - let start_pos_in_bitmap = flat_projection( - vec3u(bitmap_position_full_resolution), vec2u(4, 4), - ); - if ( - 0 == (RAY_TO_LEAF_OCCUPANCY_BITMASK_LUT[start_pos_in_bitmap][direction_lut_index * 2] - & occupancy_bitmap_lsb) - && 0 == (RAY_TO_LEAF_OCCUPANCY_BITMASK_LUT[start_pos_in_bitmap][direction_lut_index * 2 + 1] - & occupancy_bitmap_msb) - ){ - result.hit = false; - return result; - } - } - - let voxel_brick_index = u32(flat_projection( - vec3u(current_index), - vec2u(dimension, dimension) - )); - if !is_empty(voxels[brick_index_start + voxel_brick_index]) - { - result.hit = true; - result.index = vec3u(current_index); - return result; - } - - let step = dda_step_to_next_sibling( - ray, - ray_current_distance, - current_bounds, - ray_scale_factors - ); - current_bounds.min_position = current_bounds.min_position + vec3f(step); - current_index = current_index + vec3i(step); - } - return result; -} - -struct OctreeRayIntersection { - hit: bool, - albedo : vec4, - content: u32, - collision_point: vec3f, - impact_normal: vec3f, -} - -const max_depth = 20; // the depth for an octree the size of 1048576 - // which would be approximately 10 km in case 1 voxel is 1 cm -fn get_by_ray(ray_: Line) -> OctreeRayIntersection{ - var result: OctreeRayIntersection; - let dimension = octreeMetaData.voxel_brick_dim; - let voxelement_count = arrayLength(&voxels); - let node_count = arrayLength(&nodes); - - // Eliminate all zeroes within the direction of the ray - var ray = ray_; - if 0. == ray.direction.x { - ray.direction.x = FLOAT_ERROR_TOLERANCE; - } - if 0. == ray.direction.y { - ray.direction.y = FLOAT_ERROR_TOLERANCE; - } - if 0. == ray.direction.z { - ray.direction.z = FLOAT_ERROR_TOLERANCE; - } - - let ray_scale_factors = get_dda_scale_factors(ray); - let direction_lut_index = hash_direction(ray.direction); - - var root_bounds = Cube(vec3(0.,0.,0.), f32(octreeMetaData.octree_size)); - var ray_current_distance: f32 = 0.0; - var node_stack: array; - var node_stack_i: i32 = 0; - let unit_in_bitmap_space = 4. / f32(dimension); - - let root_intersection = cube_intersect_ray(root_bounds, ray); - if(root_intersection.hit){ - ray_current_distance = impact_or(root_intersection, 0.); - let target_octant = hash_region( - point_in_ray_at_distance(ray, ray_current_distance) - root_bounds.min_position, - root_bounds.size, - ); - node_stack[0] = new_node_stack_item( - root_bounds, - OCTREE_ROOT_NODE_KEY, - nodes[OCTREE_ROOT_NODE_KEY].sized_node_meta, - target_octant - ); - node_stack_i = 1; - } - while(0 < node_stack_i && node_stack_i < max_depth) { - var current_bounds = node_stack[node_stack_i - 1].bounds; - var current_node = nodes[node_stack[node_stack_i - 1].node]; //!NOTE: should be const, but then it can not be indexed dynamically - var target_octant = node_stack[node_stack_i - 1].target_octant; - - var leaf_miss = false; - if is_leaf(current_node.sized_node_meta) { - let leaf_brick_hit = traverse_brick( - ray, &ray_current_distance, current_node.voxels_start_at, - children_buffer[current_node.children_starts_at], - children_buffer[current_node.children_starts_at + 1], - current_bounds, ray_scale_factors, direction_lut_index, - unit_in_bitmap_space, dimension - ); - if leaf_brick_hit.hit == true { - let hit_in_voxels = ( - current_node.voxels_start_at - + u32(flat_projection( leaf_brick_hit.index, vec2u(dimension, dimension ))) - ); - current_bounds.size /= f32(dimension); - current_bounds.min_position = current_bounds.min_position - + vec3f(leaf_brick_hit.index) * current_bounds.size; - result.hit = true; - result.albedo = voxels[hit_in_voxels].albedo; - result.content = voxels[hit_in_voxels].content; - result.collision_point = point_in_ray_at_distance(ray, ray_current_distance); - result.impact_normal = cube_impact_normal(current_bounds, result.collision_point); - return result; - } - leaf_miss = true; - } - if( leaf_miss - || target_octant == OOB_OCTANT - || ( 0 == ( - get_node_occupancy_bitmap(current_node.sized_node_meta) - | RAY_TO_NODE_OCCUPANCY_BITMASK_LUT[target_octant][direction_lut_index] - )) - || 0 == get_node_occupancy_bitmap(current_node.sized_node_meta) - ){ - // POP - let popped_target = node_stack[node_stack_i - 1]; - node_stack_i -= 1; - if(0 < node_stack_i){ - let step_vec = dda_step_to_next_sibling( - ray, - &ray_current_distance, - popped_target.bounds, - ray_scale_factors - ); - node_stack[node_stack_i - 1].target_octant = step_octant( - node_stack[node_stack_i - 1].target_octant, - step_vec - ); - } - continue; - } - - var target_bounds = child_bounds_for(current_bounds, target_octant); - var target_child_key = children_buffer[current_node.children_starts_at + target_octant]; - let target_is_empty = ( - target_child_key >= node_count //!crate::object_pool::key_is_valid - || 0 == ( - get_node_occupancy_bitmap( current_node.sized_node_meta ) - & ( // crate::spatial::math::octant_bitmask - 0x00000001u << (target_octant & 0x000000FF) - ) - ) - ); - - if !target_is_empty { - // PUSH - let child_target_octant = hash_region( - (point_in_ray_at_distance(ray, ray_current_distance) - target_bounds.min_position), - target_bounds.size - ); - node_stack[node_stack_i] = new_node_stack_item( - target_bounds, - target_child_key, - nodes[target_child_key].sized_node_meta, - child_target_octant - ); - node_stack_i += 1; - } else { - // ADVANCE - loop { - let step_vec = dda_step_to_next_sibling( - ray, - &ray_current_distance, - target_bounds, - ray_scale_factors - ); - target_octant = step_octant(target_octant, step_vec); - if OOB_OCTANT != target_octant { - target_bounds = child_bounds_for(current_bounds, target_octant); - target_child_key = - children_buffer[current_node.children_starts_at + target_octant]; - } - - if ( - target_octant == OOB_OCTANT - || ( - target_child_key < node_count //crate::object_pool::key_is_valid - && 0 != ( - get_node_occupancy_bitmap( current_node.sized_node_meta ) - & (0x00000001u << (target_octant & 0x000000FF)) // crate::spatial::math::octant_bitmask - ) - && 0 != ( - get_node_occupancy_bitmap(nodes[target_child_key].sized_node_meta) - & RAY_TO_NODE_OCCUPANCY_BITMASK_LUT[hash_region( - point_in_ray_at_distance(ray, ray_current_distance) - target_bounds.min_position, - target_bounds.size - )][direction_lut_index] - ) - ) - ){ - node_stack[node_stack_i - 1].target_octant = target_octant; - break; - } - } - } - } - result.hit = false; - return result; -} - -struct Voxelement { - albedo : vec4, - content: u32, -} - -fn is_empty(e: Voxelement) -> bool { - return ( - 0. == e.albedo.r - && 0. == e.albedo.g - && 0. == e.albedo.b - && 0. == e.albedo.a - && 0 == e.content - ); -} - -struct SizedNode { - sized_node_meta: u32, - children_starts_at: u32, - voxels_start_at: u32, -} - -const OCTREE_ROOT_NODE_KEY = 0u; -struct OctreeMetaData { - octree_size: u32, - voxel_brick_dim: u32, - ambient_light_color: vec4f, - ambient_light_position: vec3f, -} - -struct Viewport { - origin: vec3f, - direction: vec3f, - w_h_fov: vec3f, -} - -@group(0) @binding(0) -var output_texture: texture_storage_2d; - -@group(0) @binding(1) -var output_texture_render: texture_2d; - -@group(0) @binding(2) -var output_texture_sampler: sampler; - -@group(0) @binding(3) -var viewport: Viewport; - -@group(1) @binding(0) -var octreeMetaData: OctreeMetaData; - -@group(1) @binding(1) -var nodes: array; - -@group(1) @binding(2) -var children_buffer: array; - -@group(1) @binding(3) -var voxels: array; - -@compute @workgroup_size(8, 8, 1) -fn update( - @builtin(global_invocation_id) invocation_id: vec3, - @builtin(num_workgroups) num_workgroups: vec3, -) { - let pixel_location = vec2u(invocation_id.xy); - let pixel_location_normalized = vec2f( - f32(invocation_id.x) / f32(num_workgroups.x * 8), - f32(invocation_id.y) / f32(num_workgroups.y * 8) - ); - let viewport_up_direction = vec3f(0., 1., 0.); - let viewport_right_direction = normalize(cross( - viewport_up_direction, viewport.direction - )); - let viewport_bottom_left = viewport.origin - + (viewport.direction * viewport.w_h_fov.z) - - (viewport_right_direction * (viewport.w_h_fov.x / 2.)) - - (viewport_up_direction * (viewport.w_h_fov.y / 2.)) - ; - let ray_endpoint = viewport_bottom_left - + viewport_right_direction * viewport.w_h_fov.x * f32(pixel_location_normalized.x) - + viewport_up_direction * viewport.w_h_fov.y * (1. - f32(pixel_location_normalized.y)) - ; - var ray = Line(ray_endpoint, normalize(ray_endpoint - viewport.origin)); - - var ray_result = get_by_ray(ray); - var rgb_result = vec3f(0.5,0.5,0.5); - if ray_result.hit == true { - let diffuse_light_strength = ( - dot(ray_result.impact_normal, vec3f(-0.5,0.5,-0.5)) / 2. + 0.5 - ); - let result_with_lights = ray_result.albedo.rgb * diffuse_light_strength; - rgb_result = result_with_lights.rgb; - } - - textureStore(output_texture, pixel_location, vec4f(rgb_result, 1.)); -} - -struct VertexOutput { - @builtin(position) position: vec4, - @location(0) uv: vec2, -} - -@vertex -fn vs_main(@builtin(vertex_index) in_vertex_index: u32) -> VertexOutput { - // map the input vertex index values to rectangle x,y coordinates - var x = 0.; var y = 0.; - if 0 == in_vertex_index || 3 == in_vertex_index { - x = -1.; - y = -1.; - }else if 1 == in_vertex_index { - x = -1.; - y = 1.; - }else if 2 == in_vertex_index || 4 == in_vertex_index { - x = 1.; - y = 1.; - }else if 5 == in_vertex_index { - x = 1.; - y = -1.; - } - - let pos = vec4f(x, y, 0.0, 1.0); - let uv = vec2f((x + 1.) / 2.,(y + 1.) / 2.); - return VertexOutput(pos,uv); -} - -@fragment -fn fs_main(vertex_output: VertexOutput) -> @location(0) vec4 { - // Display initial texture -/* let condition = vertex_output.uv < vec2f(0.5,0.5); - if condition.x && condition.y { - //return textureLoad(output_texture, vec2u(vertex_output.uv * 100)); - return textureSample( - output_texture_render, output_texture_sampler, - vertex_output.uv - ); - } - - return vec4f(vertex_output.uv, 0.0, 1.0); -*/ - - // Display viewport - //return vec4f(0.,0.2,1., 1.); - return vec4f(viewport.origin, 1.); -} - -// Note: should be const -var OCTANT_STEP_RESULT_LUT: array, 3>, 3> = array, 3>, 3>( - array, 3>( - array(143165576,671647880,2284357768), - array(1216874632,1749559304,2288551976), - array(2290632840,2290640968,2290649192) - ), - array, 3>( - array(277383304,839944328,2285013128), - array(1418203272,1985229328,2289469490), - array(2290635912,2290644564,2290649206) - ), - array, 3>( - array(2173208712,2206304392,2290321544), - array(2240315784,2273674113,2290583683), - array(2290648456,2290648965,2290649223) - ) -); - -// Note: should be const -var RAY_TO_NODE_OCCUPANCY_BITMASK_LUT: array, 8> = array, 8>( - array(1, 3, 5, 15, 17, 51, 85, 255), - array(3, 2, 15, 10, 51, 34, 255, 170), - array(5, 15, 4, 12, 85, 255, 68, 204), - array(15, 10, 12, 8, 255, 170, 204, 136), - array(17, 51, 85, 255, 16, 48, 80, 240), - array(51, 34, 255, 170, 48, 32, 240, 160), - array(85, 255, 68, 204, 80, 240, 64, 192), - array(255, 170, 204, 136, 240, 160, 192, 128), -); - -// Note: should be const -var RAY_TO_LEAF_OCCUPANCY_BITMASK_LUT: array, 64> = array, 64>( - array(1,0,15,0,65537,65537,983055,983055,4369,0,65535,0,286331153,286331153,4294967295,4294967295,), - array(3,0,14,0,196611,196611,917518,917518,13107,0,61166,0,858993459,858993459,4008636142,4008636142,), - array(7,0,12,0,458759,458759,786444,786444,30583,0,52428,0,2004318071,2004318071,3435973836,3435973836,), - array(15,0,8,0,983055,983055,524296,524296,65535,0,34952,0,4294967295,4294967295,2290649224,2290649224,), - array(17,0,255,0,1114129,1114129,16711935,16711935,4368,0,65520,0,286265616,286265616,4293984240,4293984240,), - array(51,0,238,0,3342387,3342387,15597806,15597806,13104,0,61152,0,858796848,858796848,4007718624,4007718624,), - array(119,0,204,0,7798903,7798903,13369548,13369548,30576,0,52416,0,2003859312,2003859312,3435187392,3435187392,), - array(255,0,136,0,16711935,16711935,8913032,8913032,65520,0,34944,0,4293984240,4293984240,2290124928,2290124928,), - array(273,0,4095,0,17891601,17891601,268374015,268374015,4352,0,65280,0,285217024,285217024,4278255360,4278255360,), - array(819,0,3822,0,53674803,53674803,250482414,250482414,13056,0,60928,0,855651072,855651072,3993038336,3993038336,), - array(1911,0,3276,0,125241207,125241207,214699212,214699212,30464,0,52224,0,1996519168,1996519168,3422604288,3422604288,), - array(4095,0,2184,0,268374015,268374015,143132808,143132808,65280,0,34816,0,4278255360,4278255360,2281736192,2281736192,), - array(4369,0,65535,0,286331153,286331153,4294967295,4294967295,4096,0,61440,0,268439552,268439552,4026593280,4026593280,), - array(13107,0,61166,0,858993459,858993459,4008636142,4008636142,12288,0,57344,0,805318656,805318656,3758153728,3758153728,), - array(30583,0,52428,0,2004318071,2004318071,3435973836,3435973836,28672,0,49152,0,1879076864,1879076864,3221274624,3221274624,), - array(65535,0,34952,0,4294967295,4294967295,2290649224,2290649224,61440,0,32768,0,4026593280,4026593280,2147516416,2147516416,), - array(65537,0,983055,0,65536,65537,983040,983055,286331153,0,4294967295,0,286326784,286331153,4294901760,4294967295,), - array(196611,0,917518,0,196608,196611,917504,917518,858993459,0,4008636142,0,858980352,858993459,4008574976,4008636142,), - array(458759,0,786444,0,458752,458759,786432,786444,2004318071,0,3435973836,0,2004287488,2004318071,3435921408,3435973836,), - array(983055,0,524296,0,983040,983055,524288,524296,4294967295,0,2290649224,0,4294901760,4294967295,2290614272,2290649224,), - array(1114129,0,16711935,0,1114112,1114129,16711680,16711935,286265616,0,4293984240,0,286261248,286265616,4293918720,4293984240,), - array(3342387,0,15597806,0,3342336,3342387,15597568,15597806,858796848,0,4007718624,0,858783744,858796848,4007657472,4007718624,), - array(7798903,0,13369548,0,7798784,7798903,13369344,13369548,2003859312,0,3435187392,0,2003828736,2003859312,3435134976,3435187392,), - array(16711935,0,8913032,0,16711680,16711935,8912896,8913032,4293984240,0,2290124928,0,4293918720,4293984240,2290089984,2290124928,), - array(17891601,0,268374015,0,17891328,17891601,268369920,268374015,285217024,0,4278255360,0,285212672,285217024,4278190080,4278255360,), - array(53674803,0,250482414,0,53673984,53674803,250478592,250482414,855651072,0,3993038336,0,855638016,855651072,3992977408,3993038336,), - array(125241207,0,214699212,0,125239296,125241207,214695936,214699212,1996519168,0,3422604288,0,1996488704,1996519168,3422552064,3422604288,), - array(268374015,0,143132808,0,268369920,268374015,143130624,143132808,4278255360,0,2281736192,0,4278190080,4278255360,2281701376,2281736192,), - array(286331153,0,4294967295,0,286326784,286331153,4294901760,4294967295,268439552,0,4026593280,0,268435456,268439552,4026531840,4026593280,), - array(858993459,0,4008636142,0,858980352,858993459,4008574976,4008636142,805318656,0,3758153728,0,805306368,805318656,3758096384,3758153728,), - array(2004318071,0,3435973836,0,2004287488,2004318071,3435921408,3435973836,1879076864,0,3221274624,0,1879048192,1879076864,3221225472,3221274624,), - array(4294967295,0,2290649224,0,4294901760,4294967295,2290614272,2290649224,4026593280,0,2147516416,0,4026531840,4026593280,2147483648,2147516416,), - array(65537,1,983055,15,0,65537,0,983055,286331153,4369,4294967295,65535,0,286331153,0,4294967295,), - array(196611,3,917518,14,0,196611,0,917518,858993459,13107,4008636142,61166,0,858993459,0,4008636142,), - array(458759,7,786444,12,0,458759,0,786444,2004318071,30583,3435973836,52428,0,2004318071,0,3435973836,), - array(983055,15,524296,8,0,983055,0,524296,4294967295,65535,2290649224,34952,0,4294967295,0,2290649224,), - array(1114129,17,16711935,255,0,1114129,0,16711935,286265616,4368,4293984240,65520,0,286265616,0,4293984240,), - array(3342387,51,15597806,238,0,3342387,0,15597806,858796848,13104,4007718624,61152,0,858796848,0,4007718624,), - array(7798903,119,13369548,204,0,7798903,0,13369548,2003859312,30576,3435187392,52416,0,2003859312,0,3435187392,), - array(16711935,255,8913032,136,0,16711935,0,8913032,4293984240,65520,2290124928,34944,0,4293984240,0,2290124928,), - array(17891601,273,268374015,4095,0,17891601,0,268374015,285217024,4352,4278255360,65280,0,285217024,0,4278255360,), - array(53674803,819,250482414,3822,0,53674803,0,250482414,855651072,13056,3993038336,60928,0,855651072,0,3993038336,), - array(125241207,1911,214699212,3276,0,125241207,0,214699212,1996519168,30464,3422604288,52224,0,1996519168,0,3422604288,), - array(268374015,4095,143132808,2184,0,268374015,0,143132808,4278255360,65280,2281736192,34816,0,4278255360,0,2281736192,), - array(286331153,4369,4294967295,65535,0,286331153,0,4294967295,268439552,4096,4026593280,61440,0,268439552,0,4026593280,), - array(858993459,13107,4008636142,61166,0,858993459,0,4008636142,805318656,12288,3758153728,57344,0,805318656,0,3758153728,), - array(2004318071,30583,3435973836,52428,0,2004318071,0,3435973836,1879076864,28672,3221274624,49152,0,1879076864,0,3221274624,), - array(4294967295,65535,2290649224,34952,0,4294967295,0,2290649224,4026593280,61440,2147516416,32768,0,4026593280,0,2147516416,), - array(65537,65537,983055,983055,0,65536,0,983040,286331153,286331153,4294967295,4294967295,0,286326784,0,4294901760,), - array(196611,196611,917518,917518,0,196608,0,917504,858993459,858993459,4008636142,4008636142,0,858980352,0,4008574976,), - array(458759,458759,786444,786444,0,458752,0,786432,2004318071,2004318071,3435973836,3435973836,0,2004287488,0,3435921408,), - array(983055,983055,524296,524296,0,983040,0,524288,4294967295,4294967295,2290649224,2290649224,0,4294901760,0,2290614272,), - array(1114129,1114129,16711935,16711935,0,1114112,0,16711680,286265616,286265616,4293984240,4293984240,0,286261248,0,4293918720,), - array(3342387,3342387,15597806,15597806,0,3342336,0,15597568,858796848,858796848,4007718624,4007718624,0,858783744,0,4007657472,), - array(7798903,7798903,13369548,13369548,0,7798784,0,13369344,2003859312,2003859312,3435187392,3435187392,0,2003828736,0,3435134976,), - array(16711935,16711935,8913032,8913032,0,16711680,0,8912896,4293984240,4293984240,2290124928,2290124928,0,4293918720,0,2290089984,), - array(17891601,17891601,268374015,268374015,0,17891328,0,268369920,285217024,285217024,4278255360,4278255360,0,285212672,0,4278190080,), - array(53674803,53674803,250482414,250482414,0,53673984,0,250478592,855651072,855651072,3993038336,3993038336,0,855638016,0,3992977408,), - array(125241207,125241207,214699212,214699212,0,125239296,0,214695936,1996519168,1996519168,3422604288,3422604288,0,1996488704,0,3422552064,), - array(268374015,268374015,143132808,143132808,0,268369920,0,143130624,4278255360,4278255360,2281736192,2281736192,0,4278190080,0,2281701376,), - array(286331153,286331153,4294967295,4294967295,0,286326784,0,4294901760,268439552,268439552,4026593280,4026593280,0,268435456,0,4026531840,), - array(858993459,858993459,4008636142,4008636142,0,858980352,0,4008574976,805318656,805318656,3758153728,3758153728,0,805306368,0,3758096384,), - array(2004318071,2004318071,3435973836,3435973836,0,2004287488,0,3435921408,1879076864,1879076864,3221274624,3221274624,0,1879048192,0,3221225472,), - array(4294967295,4294967295,2290649224,2290649224,0,4294901760,0,2290614272,4026593280,4026593280,2147516416,2147516416,0,4026531840,0,2147483648,), -); diff --git a/src/octree/raytracing/wgpu/render.rs b/src/octree/raytracing/wgpu/render.rs deleted file mode 100644 index c0e5536..0000000 --- a/src/octree/raytracing/wgpu/render.rs +++ /dev/null @@ -1,66 +0,0 @@ -pub use crate::octree::raytracing::wgpu::types::SvxRenderApp; - -impl SvxRenderApp { - pub fn execute_pipeline(&self) { - let frame = self - .surface - .as_ref() - .expect("Render Surface not available") - .get_current_texture() - .expect("Failed to acquire next swap chain texture"); - - let view = frame - .texture - .create_view(&wgpu::TextureViewDescriptor::default()); - let mut encoder = self - .device - .as_ref() - .expect("Device Encoder not found") - .create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None }); - - encoder.copy_texture_to_texture( - self.output_texture - .as_ref() - .expect("Expected Output texture") - .as_image_copy(), - self.output_texture_render - .as_ref() - .expect("Expected Output render texture") - .as_image_copy(), - self.texture_extent, - ); - - { - let mut render_pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { - label: None, - color_attachments: &[Some(wgpu::RenderPassColorAttachment { - view: &view, - resolve_target: None, //render target! - ops: wgpu::Operations { - load: wgpu::LoadOp::Clear(wgpu::Color::GREEN), - store: wgpu::StoreOp::Store, - }, - })], - depth_stencil_attachment: None, - timestamp_writes: None, - occlusion_query_set: None, - }); - render_pass.set_pipeline(&self.pipeline.as_ref().unwrap()); - render_pass.set_bind_group( - 0, - &self - .dynamic_group - .as_ref() - .expect("Expected Dynamic Bind Group"), - &[], - ); - render_pass.draw(0..6, 0..1); - } - - self.queue - .as_ref() - .expect("Render Queue not available") - .submit(Some(encoder.finish())); - frame.present(); - } -} diff --git a/src/octree/raytracing/wgpu/types.rs b/src/octree/raytracing/wgpu/types.rs deleted file mode 100644 index 2284270..0000000 --- a/src/octree/raytracing/wgpu/types.rs +++ /dev/null @@ -1,110 +0,0 @@ -use crate::octree::Albedo; -use crate::spatial::math::vector::V3cf32; -use encase::ShaderType; -use wgpu::{Adapter, BindGroup, Buffer, Device, Queue, RenderPipeline, Surface, Texture}; - -#[derive(ShaderType)] -pub(crate) struct Voxelement { - pub(crate) albedo: Albedo, - pub(crate) content: u32, -} - -#[derive(ShaderType)] -pub(crate) struct SizedNode { - /// Composite field: - /// - Byte 1: Boolean value, true in case node is a leaf - /// - In case of internal nodes: - /// - Byte 2: TBD - /// - Byte 3: TBD - /// - Byte 4: Lvl2 Occupancy bitmap - /// - In case of leaf nodes: - /// - Byte 2: TBD - /// - Byte 3: TBD - /// - Byte 4: TBD - pub(crate) sized_node_meta: u32, - - /// index of where the data about this node is found in children_buffer - /// - In case of internal nodes: - /// - 8 Index value of node children - /// - In case of leaf nodes: - /// - Byte 1-4: Occupancy bitmap MSB - /// - Byte 5-8: Occupancy bitmap LSB - /// - Byte 9-12: TBD - /// - Byte 13-16: TBD - /// - Byte 17-20: TBD - /// - Byte 21-24: TBD - /// - Byte 25-28: TBD - /// - Byte 29-32: TBD - pub(crate) children_start_at: u32, - - /// index of where the voxel values contained in the node start inside the voxels buffer, - /// or a "none_value". Should the field contain an index, the next voxel_brick_dim^3 elements - /// inside the @voxels array count as part of the voxels associated with the node - pub(crate) voxels_start_at: u32, -} - -#[derive(ShaderType)] -pub struct OctreeMetaData { - pub(crate) octree_size: u32, - pub(crate) voxel_brick_dim: u32, - pub ambient_light_color: V3cf32, - pub ambient_light_position: V3cf32, -} - -#[derive(ShaderType, PartialEq)] -pub struct Viewport { - pub origin: V3cf32, - pub direction: V3cf32, - pub w_h_fov: V3cf32, -} - -impl Default for Viewport { - fn default() -> Self { - Self { - origin: [0., 0., -1.].into(), - direction: [0., 0., 1.].into(), - w_h_fov: [1.5, 1., 45.].into(), - } - } -} - -pub struct SvxRenderApp { - //render data and parameters - pub(crate) viewport: Viewport, - pub(crate) output_width: u32, - pub(crate) output_height: u32, - pub(crate) texture_extent: wgpu::Extent3d, - - // wgpu pipeline - pub(crate) wgpu_instance: wgpu::Instance, - pub(crate) adapter: Option, - pub(crate) surface: Option>, - pub(crate) device: Option, - pub(crate) pipeline: Option, - pub(crate) queue: Option, - - //layouts, textures and buffers - pub(crate) dynamic_group: Option, - // pub(crate) tree_group: Option, - pub(crate) output_texture: Option, - pub(crate) output_texture_render: Option, - pub(crate) viewport_buffer: Option, - // pub(crate) metadata_buffer: Option, - // pub(crate) nodes_buffer: Option, - // pub(crate) children_buffer: Option, - // pub(crate) voxels_buffer: Option, -} - -#[cfg(test)] -mod types_wgpu_byte_compatibility_tests { - use super::{OctreeMetaData, SizedNode, Viewport, Voxelement}; - use encase::ShaderType; - - #[test] - fn test_wgpu_compatibility() { - Viewport::assert_uniform_compat(); - OctreeMetaData::assert_uniform_compat(); - Voxelement::assert_uniform_compat(); - SizedNode::assert_uniform_compat(); - } -} diff --git a/src/octree/tests.rs b/src/octree/tests.rs index a91abaa..154649c 100644 --- a/src/octree/tests.rs +++ b/src/octree/tests.rs @@ -1,4 +1,5 @@ #[cfg(test)] +#[cfg(feature = "bevy_wgpu")] mod types_byte_compatibility { use crate::octree::Albedo; use encase::{ShaderType, StorageBuffer}; diff --git a/src/octree/types.rs b/src/octree/types.rs index 2d1706e..7139913 100644 --- a/src/octree/types.rs +++ b/src/octree/types.rs @@ -55,42 +55,6 @@ pub trait VoxelData { fn clear(&mut self); } -impl VoxelData for Albedo { - fn new(color: Albedo, _user_data: u32) -> Self { - color - } - - fn albedo(&self) -> Albedo { - *self - } - - fn user_data(&self) -> u32 { - 0u32 - } - - fn clear(&mut self) { - self.r = 0.; - self.r = 0.; - self.b = 0.; - self.a = 0.; - } -} - -impl From for Albedo { - fn from(value: u32) -> Self { - let a = (value & 0x000000FF) as u8; - let b = ((value & 0x0000FF00) >> 8) as u8; - let g = ((value & 0x00FF0000) >> 16) as u8; - let r = ((value & 0xFF000000) >> 24) as u8; - - Albedo::default() - .with_red(r as f32 / 255.) - .with_green(g as f32 / 255.) - .with_blue(b as f32 / 255.) - .with_alpha(a as f32 / 255.) - } -} - /// Sparse Octree of Nodes, where each node contains a brick of voxels. /// A Brick is a 3 dimensional matrix, each element of it containing a voxel. /// A Brick can be indexed directly, as opposed to the octree which is essentially a @@ -103,7 +67,7 @@ pub struct Octree { pub(in crate::octree) node_children: Vec>, // Children index values of each Node } -#[cfg_attr(feature = "wgpu", derive(encase::ShaderType))] +#[cfg_attr(feature = "bevy_wgpu", derive(encase::ShaderType))] #[derive(Default, Clone, Copy, Debug, PartialEq)] pub struct Albedo { pub r: f32, @@ -137,3 +101,39 @@ impl Albedo { self.a == 0.0 } } + +impl From for Albedo { + fn from(value: u32) -> Self { + let a = (value & 0x000000FF) as u8; + let b = ((value & 0x0000FF00) >> 8) as u8; + let g = ((value & 0x00FF0000) >> 16) as u8; + let r = ((value & 0xFF000000) >> 24) as u8; + + Albedo::default() + .with_red(r as f32 / 255.) + .with_green(g as f32 / 255.) + .with_blue(b as f32 / 255.) + .with_alpha(a as f32 / 255.) + } +} + +impl VoxelData for Albedo { + fn new(color: Albedo, _user_data: u32) -> Self { + color + } + + fn albedo(&self) -> Albedo { + *self + } + + fn user_data(&self) -> u32 { + 0u32 + } + + fn clear(&mut self) { + self.r = 0.; + self.r = 0.; + self.b = 0.; + self.a = 0.; + } +} diff --git a/src/spatial/math/tests.rs b/src/spatial/math/tests.rs index 370c330..e3c9767 100644 --- a/src/spatial/math/tests.rs +++ b/src/spatial/math/tests.rs @@ -67,7 +67,7 @@ mod intersection_tests { } #[cfg(test)] -#[cfg(feature = "wgpu")] +#[cfg(feature = "bevy_wgpu")] mod wgpu_tests { use crate::spatial::math::vector::V3cf32; use encase::StorageBuffer; diff --git a/src/spatial/math/vector.rs b/src/spatial/math/vector.rs index 41f57ee..7758015 100644 --- a/src/spatial/math/vector.rs +++ b/src/spatial/math/vector.rs @@ -264,20 +264,20 @@ impl From> for V3c { } } -#[cfg(feature = "wgpu")] +#[cfg(feature = "bevy_wgpu")] use encase::{impl_vector, vector::AsMutVectorParts, vector::AsRefVectorParts}; -#[cfg(feature = "wgpu")] +#[cfg(feature = "bevy_wgpu")] impl_vector!(3, V3cf32, f32; using From); -#[cfg(feature = "wgpu")] +#[cfg(feature = "bevy_wgpu")] impl AsRefVectorParts for V3cf32 { fn as_ref_parts(&self) -> &[f32; 3] { unsafe { &*(self as *const V3cf32 as *const [f32; 3]) } } } -#[cfg(feature = "wgpu")] +#[cfg(feature = "bevy_wgpu")] impl AsMutVectorParts for V3cf32 { fn as_mut_parts(&mut self) -> &mut [f32; 3] { unsafe { &mut *(self as *mut V3cf32 as *mut [f32; 3]) } From b463bd65e58992ca55abd3cf9e0ed69c214c69b4 Mon Sep 17 00:00:00 2001 From: davids91 Date: Fri, 12 Jul 2024 17:49:42 +0200 Subject: [PATCH 3/5] updated bevy to latest version --- Cargo.toml | 5 +-- assets/shaders/viewport_render.wgsl | 13 +++--- examples/bevy_wgpu_render.rs | 25 ++++++----- src/octree/raytracing/bevy/data.rs | 23 +++------- src/octree/raytracing/bevy/mod.rs | 7 +-- src/octree/raytracing/bevy/types.rs | 66 ++++++++++++++++++++++++----- src/octree/tests.rs | 9 +++- src/octree/types.rs | 12 +++++- src/spatial/math/tests.rs | 2 +- src/spatial/math/vector.rs | 4 +- 10 files changed, 111 insertions(+), 55 deletions(-) diff --git a/Cargo.toml b/Cargo.toml index 9342fc5..257d326 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -9,19 +9,18 @@ license = "MIT OR Apache-2.0" default = [] raytracing = ["dep:image", "dep:show-image"] serialization = ["dep:serde"] -bevy_wgpu = ["raytracing", "dep:bevy", "dep:encase"] +bevy_wgpu = ["raytracing", "dep:bevy"] [dependencies] serde = { version = "1.0.183", features = ["derive"], optional = true } bendy = { git = "https://github.com/davids91/bendy.git" , features = ["std", "serde"]} array-init = "2.1.0" -encase = { version = "0.9.0", optional = true } # for example cpu_render image = { version = "0.25.1", optional = true } show-image = { version = "0.14.0", optional = true } # for example bevy_wgpu -bevy = { version = "0.13.2", features = ["dynamic_linking"], optional = true} +bevy = { version = "0.14.0", features = [], optional = true} # debugging #linker = "/usr/bin/clang" diff --git a/assets/shaders/viewport_render.wgsl b/assets/shaders/viewport_render.wgsl index 9ae8726..8c46cb2 100644 --- a/assets/shaders/viewport_render.wgsl +++ b/assets/shaders/viewport_render.wgsl @@ -628,8 +628,7 @@ struct OctreeMetaData { struct Viewport { origin: vec3f, direction: vec3f, - size: vec2f, - fov: f32, + w_h_fov: vec3f, } @group(0) @binding(0) @@ -665,13 +664,13 @@ fn update( viewport_up_direction, viewport.direction )); let viewport_bottom_left = viewport.origin - + (viewport.direction * viewport.fov) - - (viewport_right_direction * (viewport.size.x / 2.)) - - (viewport_up_direction * (viewport.size.y / 2.)) + + (viewport.direction * viewport.w_h_fov.z) + - (viewport_right_direction * (viewport.w_h_fov.x / 2.)) + - (viewport_up_direction * (viewport.w_h_fov.y / 2.)) ; let ray_endpoint = viewport_bottom_left - + viewport_right_direction * viewport.size.x * f32(pixel_location_normalized.x) - + viewport_up_direction * viewport.size.y * (1. - f32(pixel_location_normalized.y)) + + viewport_right_direction * viewport.w_h_fov.x * f32(pixel_location_normalized.x) + + viewport_up_direction * viewport.w_h_fov.y * (1. - f32(pixel_location_normalized.y)) ; var ray = Line(ray_endpoint, normalize(ray_endpoint - viewport.origin)); diff --git a/examples/bevy_wgpu_render.rs b/examples/bevy_wgpu_render.rs index 3c3de25..dbe7850 100644 --- a/examples/bevy_wgpu_render.rs +++ b/examples/bevy_wgpu_render.rs @@ -33,7 +33,7 @@ fn main() { fn setup(mut commands: Commands, images: ResMut>) { use shocovox_rs::octree::raytracing::bevy::create_viewing_glass; - let origin = Vec3::new( + let origin = V3c::new( ARRAY_DIMENSION as f32 * 2., ARRAY_DIMENSION as f32 / 2., ARRAY_DIMENSION as f32 * -2., @@ -87,10 +87,9 @@ fn setup(mut commands: Commands, images: ResMut>) { let render_data = tree.create_bevy_view(); let viewing_glass = create_viewing_glass( &Viewport { - direction: (Vec3::new(0., 0., 0.) - origin).normalize(), origin, - size: Vec2::new(10., 10.), - fov: 3., + direction: (V3c::new(0., 0., 0.) - origin).normalized(), + w_h_fov: V3c::new(10., 10., 3.), }, DISPLAY_RESOLUTION, images, @@ -121,7 +120,7 @@ fn rotate_camera( mut viewing_glass: ResMut, ) { let angle = { - let addition = ARRAY_DIMENSION as f32 / 100000.; + let addition = ARRAY_DIMENSION as f32 / 10000.; let angle = angles_query.single().yaw + addition; if angle < 360. { angle @@ -132,26 +131,30 @@ fn rotate_camera( angles_query.single_mut().yaw = angle; let radius = ARRAY_DIMENSION as f32 * 1.3; - viewing_glass.viewport.origin = Vec3::new( + viewing_glass.viewport.origin = V3c::new( ARRAY_DIMENSION as f32 / 2. + angle.sin() * radius, - ARRAY_DIMENSION as f32 / 2., + ARRAY_DIMENSION as f32 + angle.cos() * angle.sin() * radius / 2., ARRAY_DIMENSION as f32 / 2. + angle.cos() * radius, ); - viewing_glass.viewport.direction = (Vec3::new( + viewing_glass.viewport.direction = (V3c::new( ARRAY_DIMENSION as f32 / 2., ARRAY_DIMENSION as f32 / 2., ARRAY_DIMENSION as f32 / 2., ) - viewing_glass.viewport.origin) - .normalize(); + .normalized(); } #[cfg(feature = "bevy_wgpu")] fn handle_zoom(keys: Res>, mut viewing_glass: ResMut) { if keys.pressed(KeyCode::ArrowUp) { - viewing_glass.viewport.size *= 1.1; + //viewing_glass.viewport.w_h_fov *= 1.1; + viewing_glass.viewport.w_h_fov.x *= 1.1; + viewing_glass.viewport.w_h_fov.y *= 1.1; } if keys.pressed(KeyCode::ArrowDown) { - viewing_glass.viewport.size *= 0.9; + //viewing_glass.viewport.w_h_fov *= 0.9; + viewing_glass.viewport.w_h_fov.x *= 0.9; + viewing_glass.viewport.w_h_fov.y *= 0.9; } } diff --git a/src/octree/raytracing/bevy/data.rs b/src/octree/raytracing/bevy/data.rs index 81edb7a..0004d7f 100644 --- a/src/octree/raytracing/bevy/data.rs +++ b/src/octree/raytracing/bevy/data.rs @@ -3,11 +3,9 @@ use crate::object_pool::empty_marker; use crate::octree::{ raytracing::bevy::types::{OctreeMetaData, ShocoVoxRenderData, SizedNode, Voxelement}, types::{NodeChildrenArray, NodeContent}, - Octree, VoxelData, + Octree, V3c, VoxelData, }; -use bevy::{math::Vec3, render::color::Color}; - impl Octree where T: Default + Clone + VoxelData, @@ -48,8 +46,8 @@ where let meta = OctreeMetaData { octree_size: self.octree_size, voxel_brick_dim: DIM as u32, - ambient_light_color: Color::rgba(1., 1., 1., 1.), - ambient_light_position: Vec3::new( + 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, @@ -84,17 +82,10 @@ where for z in 0..DIM { for y in 0..DIM { for x in 0..DIM { - let albedo = data[x][y][z].albedo(); - let content = data[x][y][z].user_data(); - voxels.push(Voxelement { - albedo: Color::rgba( - albedo.r as f32 / 255., - albedo.g as f32 / 255., - albedo.b as f32 / 255., - albedo.a as f32 / 255., - ), - content, - }) + voxels.push(Voxelement::new( + data[x][y][z].albedo(), + data[x][y][z].user_data(), + )) } } } diff --git a/src/octree/raytracing/bevy/mod.rs b/src/octree/raytracing/bevy/mod.rs index d4d75db..cbdb072 100644 --- a/src/octree/raytracing/bevy/mod.rs +++ b/src/octree/raytracing/bevy/mod.rs @@ -17,6 +17,7 @@ use bevy::{ prelude::IntoSystemConfigs, render::{ extract_resource::ExtractResourcePlugin, + prelude::Image, render_asset::{RenderAssetUsages, RenderAssets}, render_graph, render_graph::RenderGraph, @@ -25,7 +26,7 @@ use bevy::{ Extent3d, PipelineCache, TextureDimension, TextureFormat, TextureUsages, }, renderer::{RenderContext, RenderDevice}, - texture::{FallbackImage, Image}, + texture::{FallbackImage, GpuImage}, Render, RenderApp, RenderSet, }, }; @@ -77,7 +78,7 @@ impl FromWorld for ShocoVoxRenderPipeline { } fn prepare_bind_groups( - gpu_images: Res>, + gpu_images: Res>, fallback_image: Res, render_device: Res, mut pipeline: ResMut, @@ -140,7 +141,7 @@ impl Plugin for ShocoVoxRenderPlugin { prepare_bind_groups.in_set(RenderSet::PrepareBindGroups), ); - let mut render_graph = render_app.world.resource_mut::(); + let mut render_graph = render_app.world_mut().resource_mut::(); render_graph.add_node( ShocoVoxLabel, ShocoVoxRenderNode { diff --git a/src/octree/raytracing/bevy/types.rs b/src/octree/raytracing/bevy/types.rs index cc9a23b..c1f890d 100644 --- a/src/octree/raytracing/bevy/types.rs +++ b/src/octree/raytracing/bevy/types.rs @@ -1,23 +1,33 @@ +use crate::octree::{Albedo, V3cf32, VoxelData}; use bevy::{ asset::Handle, ecs::system::Resource, - math::{Vec2, Vec3}, reflect::TypePath, render::{ - color::Color, extract_resource::ExtractResource, + prelude::Image, render_graph::RenderLabel, render_resource::{ AsBindGroup, BindGroup, BindGroupLayout, CachedComputePipelineId, ShaderType, }, - texture::Image, }, }; #[derive(Clone, ShaderType)] pub(crate) struct Voxelement { - pub(crate) albedo: Color, + pub(crate) albedo: Albedo, pub(crate) content: u32, + _padding: V3cf32, +} + +impl Voxelement { + pub fn new(albedo: Albedo, content: u32) -> Self { + Voxelement { + albedo, + content, + _padding: V3cf32::new(0., 0., 0.), + } + } } #[derive(Clone, ShaderType)] @@ -54,20 +64,21 @@ pub(crate) struct SizedNode { pub(crate) voxels_start_at: u32, } +#[repr(C)] #[derive(Clone, ShaderType)] pub struct OctreeMetaData { - pub(crate) octree_size: u32, + pub ambient_light_color: V3cf32, pub(crate) voxel_brick_dim: u32, - pub ambient_light_color: Color, - pub ambient_light_position: Vec3, + pub ambient_light_position: V3cf32, + pub(crate) octree_size: u32, } +#[repr(C)] #[derive(Clone, Copy, ShaderType)] pub struct Viewport { - pub origin: Vec3, - pub direction: Vec3, - pub size: Vec2, - pub fov: f32, + pub origin: V3cf32, + pub direction: V3cf32, + pub w_h_fov: V3cf32, } pub struct ShocoVoxRenderPlugin { @@ -117,3 +128,36 @@ pub(crate) struct ShocoVoxRenderNode { pub(crate) ready: bool, pub(crate) resolution: [u32; 2], } + +#[cfg(test)] +mod types_wgpu_byte_compatibility_tests { + use super::{OctreeMetaData, SizedNode, Viewport, Voxelement}; + use bevy::render::render_resource::{encase::ShaderType, ShaderSize}; + + #[test] + fn test_wgpu_compatibility() { + Viewport::assert_uniform_compat(); + assert_eq!( + std::mem::size_of::(), + Viewport::SHADER_SIZE.get() as usize + ); + + OctreeMetaData::assert_uniform_compat(); + assert_eq!( + std::mem::size_of::(), + OctreeMetaData::SHADER_SIZE.get() as usize + ); + + Voxelement::assert_uniform_compat(); + assert_eq!( + std::mem::size_of::(), + Voxelement::SHADER_SIZE.get() as usize + ); + + SizedNode::assert_uniform_compat(); + assert_eq!( + std::mem::size_of::(), + SizedNode::SHADER_SIZE.get() as usize + ); + } +} diff --git a/src/octree/tests.rs b/src/octree/tests.rs index 154649c..ccacc79 100644 --- a/src/octree/tests.rs +++ b/src/octree/tests.rs @@ -2,7 +2,10 @@ #[cfg(feature = "bevy_wgpu")] mod types_byte_compatibility { use crate::octree::Albedo; - use encase::{ShaderType, StorageBuffer}; + use bevy::render::render_resource::{ + encase::{ShaderSize, StorageBuffer}, + ShaderType, + }; #[test] fn albedo_size_is_as_expected() { @@ -18,6 +21,10 @@ mod types_byte_compatibility { #[test] fn test_wgpu_compatibility() { Albedo::assert_uniform_compat(); + assert_eq!( + std::mem::size_of::(), + Albedo::SHADER_SIZE.get() as usize + ); } #[test] diff --git a/src/octree/types.rs b/src/octree/types.rs index 7139913..2aa8392 100644 --- a/src/octree/types.rs +++ b/src/octree/types.rs @@ -3,6 +3,9 @@ use crate::object_pool::ObjectPool; #[cfg(feature = "serialization")] use serde::{Deserialize, Serialize}; +#[cfg(feature = "bevy_wgpu")] +use bevy::render::render_resource::ShaderType; + #[derive(Default, Clone)] #[cfg_attr(feature = "serialization", derive(Serialize, Deserialize))] pub(crate) enum NodeContent { @@ -67,7 +70,7 @@ pub struct Octree { pub(in crate::octree) node_children: Vec>, // Children index values of each Node } -#[cfg_attr(feature = "bevy_wgpu", derive(encase::ShaderType))] +#[cfg_attr(feature = "bevy_wgpu", derive(ShaderType))] #[derive(Default, Clone, Copy, Debug, PartialEq)] pub struct Albedo { pub r: f32, @@ -77,6 +80,13 @@ pub struct Albedo { } impl Albedo { + pub const WHITE: Albedo = Albedo { + r: 1., + g: 1., + b: 1., + a: 1., + }; + pub fn with_red(mut self, r: f32) -> Self { self.r = r; self diff --git a/src/spatial/math/tests.rs b/src/spatial/math/tests.rs index e3c9767..620be90 100644 --- a/src/spatial/math/tests.rs +++ b/src/spatial/math/tests.rs @@ -70,7 +70,7 @@ mod intersection_tests { #[cfg(feature = "bevy_wgpu")] mod wgpu_tests { use crate::spatial::math::vector::V3cf32; - use encase::StorageBuffer; + use bevy::render::render_resource::encase::StorageBuffer; #[test] fn test_buffer_readback() { diff --git a/src/spatial/math/vector.rs b/src/spatial/math/vector.rs index 7758015..66bc952 100644 --- a/src/spatial/math/vector.rs +++ b/src/spatial/math/vector.rs @@ -265,7 +265,9 @@ impl From> for V3c { } #[cfg(feature = "bevy_wgpu")] -use encase::{impl_vector, vector::AsMutVectorParts, vector::AsRefVectorParts}; +use bevy::render::render_resource::encase::{ + impl_vector, vector::AsMutVectorParts, vector::AsRefVectorParts, +}; #[cfg(feature = "bevy_wgpu")] impl_vector!(3, V3cf32, f32; using From); From a6272c1162809b729c8d87cd6e2343e8114a997d Mon Sep 17 00:00:00 2001 From: davids91 Date: Mon, 15 Jul 2024 17:48:33 +0200 Subject: [PATCH 4/5] Restored Albedo to u8 components, minor fixes, clarifications --- Cargo.toml | 2 +- assets/shaders/viewport_render.wgsl | 6 +- examples/bevy_wgpu_render.rs | 17 +++-- examples/cpu_render.rs | 8 +-- src/octree/bytecode.rs | 24 +++---- src/octree/raytracing/bevy/data.rs | 16 +++-- src/octree/raytracing/bevy/types.rs | 39 ++-------- src/octree/tests.rs | 4 -- src/octree/types.rs | 108 ++++++++++++++-------------- 9 files changed, 99 insertions(+), 125 deletions(-) diff --git a/Cargo.toml b/Cargo.toml index 257d326..4962d2b 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -1,6 +1,6 @@ [package] name = "shocovox-rs" -version = "0.3.1" +version = "0.3.2" 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 8c46cb2..1ffef7c 100644 --- a/assets/shaders/viewport_render.wgsl +++ b/assets/shaders/viewport_render.wgsl @@ -597,7 +597,7 @@ fn get_by_ray(ray_: Line) -> OctreeRayIntersection{ } struct Voxelement { - albedo : vec4, + albedo : vec4f, content: u32, } @@ -619,10 +619,10 @@ struct SizedNode { const OCTREE_ROOT_NODE_KEY = 0u; struct OctreeMetaData { + ambient_light_color: vec3f, + ambient_light_position: vec3f, octree_size: u32, voxel_brick_dim: u32, - ambient_light_color: vec4f, - ambient_light_position: vec3f, } struct Viewport { diff --git a/examples/bevy_wgpu_render.rs b/examples/bevy_wgpu_render.rs index dbe7850..5ab9e1c 100644 --- a/examples/bevy_wgpu_render.rs +++ b/examples/bevy_wgpu_render.rs @@ -60,23 +60,26 @@ fn setup(mut commands: Commands, images: ResMut>) { && (ARRAY_DIMENSION / 2) <= z) { let r = if 0 == x % (ARRAY_DIMENSION / 4) { - x as f32 / ARRAY_DIMENSION as f32 + (x as f32 / ARRAY_DIMENSION as f32 * 255.) as u32 } else { - 0.5 + 128 }; let g = if 0 == y % (ARRAY_DIMENSION / 4) { - y as f32 / ARRAY_DIMENSION as f32 + (y as f32 / ARRAY_DIMENSION as f32 * 255.) as u32 } else { - 0.5 + 128 }; let b = if 0 == z % (ARRAY_DIMENSION / 4) { - z as f32 / ARRAY_DIMENSION as f32 + (z as f32 / ARRAY_DIMENSION as f32 * 255.) as u32 } else { - 0.5 + 128 }; tree.insert( &V3c::new(x, y, z), - Albedo::default().with_red(r).with_green(g).with_blue(b), + Albedo::default() + .with_red(r as u8) + .with_green(g as u8) + .with_blue(b as u8), ) .ok() .unwrap(); diff --git a/examples/cpu_render.rs b/examples/cpu_render.rs index 482599f..e90ab48 100644 --- a/examples/cpu_render.rs +++ b/examples/cpu_render.rs @@ -30,10 +30,10 @@ fn main() { tree.insert( &V3c::new(x, y, z), Albedo::default() - .with_red(x as f32 / TREE_SIZE as f32) - .with_green(y as f32 / TREE_SIZE as f32) - .with_blue(z as f32 / TREE_SIZE as f32) - .with_alpha(1.), + .with_red((255 as f32 * x as f32 / TREE_SIZE as f32) as u8) + .with_green((255 as f32 * y as f32 / TREE_SIZE as f32) as u8) + .with_blue((255 as f32 * z as f32 / TREE_SIZE as f32) as u8) + .with_alpha(255), ) .ok() .unwrap(); diff --git a/src/octree/bytecode.rs b/src/octree/bytecode.rs index 562a370..8aded48 100644 --- a/src/octree/bytecode.rs +++ b/src/octree/bytecode.rs @@ -8,37 +8,37 @@ use bendy::{ impl<'obj, 'ser, T: Clone + VoxelData, const DIM: usize> NodeContent { fn encode_single(data: &T, encoder: &mut Encoder) -> Result<(), BencodeError> { let color = data.albedo(); - encoder.emit((color.r * 1000.) as i32)?; - encoder.emit((color.g * 1000.) as i32)?; - encoder.emit((color.b * 1000.) as i32)?; - encoder.emit((color.a * 1000.) as i32)?; + encoder.emit(color.r)?; + encoder.emit(color.g)?; + encoder.emit(color.b)?; + encoder.emit(color.a)?; encoder.emit(data.user_data()) } fn decode_single(list: &mut ListDecoder<'obj, 'ser>) -> Result { let r = match list.next_object()?.unwrap() { - Object::Integer(i) => Ok(i.parse::().ok().unwrap()), + Object::Integer(i) => Ok(i.parse::().ok().unwrap()), _ => Err(bendy::decoding::Error::unexpected_token( "int field red color component", "Something else", )), }?; let g = match list.next_object()?.unwrap() { - Object::Integer(i) => Ok(i.parse::().ok().unwrap()), + Object::Integer(i) => Ok(i.parse::().ok().unwrap()), _ => Err(bendy::decoding::Error::unexpected_token( "int field green color component", "Something else", )), }?; let b = match list.next_object()?.unwrap() { - Object::Integer(i) => Ok(i.parse::().ok().unwrap()), + Object::Integer(i) => Ok(i.parse::().ok().unwrap()), _ => Err(bendy::decoding::Error::unexpected_token( "int field blue color component", "Something else", )), }?; let a = match list.next_object()?.unwrap() { - Object::Integer(i) => Ok(i.parse::().ok().unwrap()), + Object::Integer(i) => Ok(i.parse::().ok().unwrap()), _ => Err(bendy::decoding::Error::unexpected_token( "int field alpha color component", "Something else", @@ -49,10 +49,10 @@ impl<'obj, 'ser, T: Clone + VoxelData, const DIM: usize> NodeContent { _ => 0, }; let albedo = Albedo::default() - .with_red(r as f32 / 1000.) - .with_green(g as f32 / 1000.) - .with_blue(b as f32 / 1000.) - .with_alpha(a as f32 / 1000.); + .with_red(r) + .with_green(g) + .with_blue(b) + .with_alpha(a); Ok(VoxelData::new(albedo, user_data)) } } diff --git a/src/octree/raytracing/bevy/data.rs b/src/octree/raytracing/bevy/data.rs index 0004d7f..34d2dd7 100644 --- a/src/octree/raytracing/bevy/data.rs +++ b/src/octree/raytracing/bevy/data.rs @@ -1,10 +1,10 @@ use crate::object_pool::empty_marker; - use crate::octree::{ raytracing::bevy::types::{OctreeMetaData, ShocoVoxRenderData, SizedNode, Voxelement}, types::{NodeChildrenArray, NodeContent}, Octree, V3c, VoxelData, }; +use bevy::math::Vec4; impl Octree where @@ -82,10 +82,16 @@ where for z in 0..DIM { for y in 0..DIM { for x in 0..DIM { - voxels.push(Voxelement::new( - data[x][y][z].albedo(), - data[x][y][z].user_data(), - )) + let albedo = data[x][y][z].albedo(); + voxels.push(Voxelement { + albedo: Vec4::new( + albedo.r as f32 / 255., + albedo.g as f32 / 255., + albedo.b as f32 / 255., + albedo.a as f32 / 255., + ), + content: data[x][y][z].user_data(), + }) } } } diff --git a/src/octree/raytracing/bevy/types.rs b/src/octree/raytracing/bevy/types.rs index c1f890d..e8a450c 100644 --- a/src/octree/raytracing/bevy/types.rs +++ b/src/octree/raytracing/bevy/types.rs @@ -1,7 +1,8 @@ -use crate::octree::{Albedo, V3cf32, VoxelData}; +use crate::octree::V3cf32; use bevy::{ asset::Handle, ecs::system::Resource, + math::Vec4, reflect::TypePath, render::{ extract_resource::ExtractResource, @@ -15,19 +16,8 @@ use bevy::{ #[derive(Clone, ShaderType)] pub(crate) struct Voxelement { - pub(crate) albedo: Albedo, + pub(crate) albedo: Vec4, pub(crate) content: u32, - _padding: V3cf32, -} - -impl Voxelement { - pub fn new(albedo: Albedo, content: u32) -> Self { - Voxelement { - albedo, - content, - _padding: V3cf32::new(0., 0., 0.), - } - } } #[derive(Clone, ShaderType)] @@ -64,16 +54,14 @@ pub(crate) struct SizedNode { pub(crate) voxels_start_at: u32, } -#[repr(C)] #[derive(Clone, ShaderType)] pub struct OctreeMetaData { pub ambient_light_color: V3cf32, - pub(crate) voxel_brick_dim: u32, pub ambient_light_position: V3cf32, pub(crate) octree_size: u32, + pub(crate) voxel_brick_dim: u32, } -#[repr(C)] #[derive(Clone, Copy, ShaderType)] pub struct Viewport { pub origin: V3cf32, @@ -137,27 +125,8 @@ mod types_wgpu_byte_compatibility_tests { #[test] fn test_wgpu_compatibility() { Viewport::assert_uniform_compat(); - assert_eq!( - std::mem::size_of::(), - Viewport::SHADER_SIZE.get() as usize - ); - OctreeMetaData::assert_uniform_compat(); - assert_eq!( - std::mem::size_of::(), - OctreeMetaData::SHADER_SIZE.get() as usize - ); - Voxelement::assert_uniform_compat(); - assert_eq!( - std::mem::size_of::(), - Voxelement::SHADER_SIZE.get() as usize - ); - SizedNode::assert_uniform_compat(); - assert_eq!( - std::mem::size_of::(), - SizedNode::SHADER_SIZE.get() as usize - ); } } diff --git a/src/octree/tests.rs b/src/octree/tests.rs index ccacc79..50f3a66 100644 --- a/src/octree/tests.rs +++ b/src/octree/tests.rs @@ -21,10 +21,6 @@ mod types_byte_compatibility { #[test] fn test_wgpu_compatibility() { Albedo::assert_uniform_compat(); - assert_eq!( - std::mem::size_of::(), - Albedo::SHADER_SIZE.get() as usize - ); } #[test] diff --git a/src/octree/types.rs b/src/octree/types.rs index 2aa8392..f16b672 100644 --- a/src/octree/types.rs +++ b/src/octree/types.rs @@ -3,9 +3,6 @@ use crate::object_pool::ObjectPool; #[cfg(feature = "serialization")] use serde::{Deserialize, Serialize}; -#[cfg(feature = "bevy_wgpu")] -use bevy::render::render_resource::ShaderType; - #[derive(Default, Clone)] #[cfg_attr(feature = "serialization", derive(Serialize, Deserialize))] pub(crate) enum NodeContent { @@ -58,6 +55,42 @@ pub trait VoxelData { fn clear(&mut self); } +impl VoxelData for Albedo { + fn new(color: Albedo, _user_data: u32) -> Self { + color + } + + fn albedo(&self) -> Albedo { + *self + } + + fn user_data(&self) -> u32 { + 0u32 + } + + fn clear(&mut self) { + self.r = 0; + self.r = 0; + self.b = 0; + self.a = 0; + } +} + +impl From for Albedo { + fn from(value: u32) -> Self { + let a = (value & 0x000000FF) as u8; + let b = ((value & 0x0000FF00) >> 8) as u8; + let g = ((value & 0x00FF0000) >> 16) as u8; + let r = ((value & 0xFF000000) >> 24) as u8; + + Albedo::default() + .with_red(r) + .with_green(g) + .with_blue(b) + .with_alpha(a) + } +} + /// Sparse Octree of Nodes, where each node contains a brick of voxels. /// A Brick is a 3 dimensional matrix, each element of it containing a voxel. /// A Brick can be indexed directly, as opposed to the octree which is essentially a @@ -70,80 +103,47 @@ pub struct Octree { pub(in crate::octree) node_children: Vec>, // Children index values of each Node } -#[cfg_attr(feature = "bevy_wgpu", derive(ShaderType))] #[derive(Default, Clone, Copy, Debug, PartialEq)] pub struct Albedo { - pub r: f32, - pub g: f32, - pub b: f32, - pub a: f32, + pub r: u8, + pub g: u8, + pub b: u8, + pub a: u8, } impl Albedo { - pub const WHITE: Albedo = Albedo { - r: 1., - g: 1., - b: 1., - a: 1., - }; - - pub fn with_red(mut self, r: f32) -> Self { + pub fn with_red(mut self, r: u8) -> Self { self.r = r; self } - pub fn with_green(mut self, g: f32) -> Self { + pub fn with_green(mut self, g: u8) -> Self { self.g = g; self } - pub fn with_blue(mut self, b: f32) -> Self { + pub fn with_blue(mut self, b: u8) -> Self { self.b = b; self } - pub fn with_alpha(mut self, a: f32) -> Self { + pub fn with_alpha(mut self, a: u8) -> Self { self.a = a; self } pub fn is_transparent(&self) -> bool { - self.a == 0.0 - } -} - -impl From for Albedo { - fn from(value: u32) -> Self { - let a = (value & 0x000000FF) as u8; - let b = ((value & 0x0000FF00) >> 8) as u8; - let g = ((value & 0x00FF0000) >> 16) as u8; - let r = ((value & 0xFF000000) >> 24) as u8; - - Albedo::default() - .with_red(r as f32 / 255.) - .with_green(g as f32 / 255.) - .with_blue(b as f32 / 255.) - .with_alpha(a as f32 / 255.) + self.a == 0 } } -impl VoxelData for Albedo { - fn new(color: Albedo, _user_data: u32) -> Self { - color - } - - fn albedo(&self) -> Albedo { - *self - } - - fn user_data(&self) -> u32 { - 0u32 - } - - fn clear(&mut self) { - self.r = 0.; - self.r = 0.; - self.b = 0.; - self.a = 0.; - } +#[test] +fn albedo_size_is_4_bytes() { + const SIZE: usize = std::mem::size_of::(); + const EXPECTED_SIZE: usize = 4; + assert_eq!( + SIZE, EXPECTED_SIZE, + "RGBA should be {} bytes wide but was {}", + EXPECTED_SIZE, SIZE + ); } From a77fd8a387fc25a0b0f2ce2b8301b313cc8b9e44 Mon Sep 17 00:00:00 2001 From: davids91 Date: Fri, 19 Jul 2024 18:53:31 +0200 Subject: [PATCH 5/5] Eliminated Stackoverflow on bigger brick dimensions --- Cargo.toml | 1 - examples/bevy_wgpu_render.rs | 2 +- src/octree/bytecode.rs | 15 ++++++--------- src/octree/detail.rs | 6 ++---- src/octree/mod.rs | 5 ++++- src/octree/raytracing/raytracing_on_cpu.rs | 5 +++-- src/octree/update.rs | 5 ++++- 7 files changed, 20 insertions(+), 19 deletions(-) diff --git a/Cargo.toml b/Cargo.toml index 4962d2b..e18609e 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -14,7 +14,6 @@ bevy_wgpu = ["raytracing", "dep:bevy"] [dependencies] serde = { version = "1.0.183", features = ["derive"], optional = true } bendy = { git = "https://github.com/davids91/bendy.git" , features = ["std", "serde"]} -array-init = "2.1.0" # for example cpu_render image = { version = "0.25.1", optional = true } show-image = { version = "0.14.0", optional = true } diff --git a/examples/bevy_wgpu_render.rs b/examples/bevy_wgpu_render.rs index 5ab9e1c..2979695 100644 --- a/examples/bevy_wgpu_render.rs +++ b/examples/bevy_wgpu_render.rs @@ -41,7 +41,7 @@ fn setup(mut commands: Commands, images: ResMut>) { commands.spawn(DomePosition { yaw: 0. }); // fill octree with data - let mut tree = shocovox_rs::octree::Octree::::new(ARRAY_DIMENSION) + let mut tree = shocovox_rs::octree::Octree::::new(ARRAY_DIMENSION) .ok() .unwrap(); diff --git a/src/octree/bytecode.rs b/src/octree/bytecode.rs index 8aded48..b7309ee 100644 --- a/src/octree/bytecode.rs +++ b/src/octree/bytecode.rs @@ -89,7 +89,7 @@ use bendy::decoding::{FromBencode, Object}; use super::types::Albedo; impl FromBencode for NodeContent where - T: PartialEq + Default + Clone + VoxelData, + T: PartialEq + Default + Clone + Copy + VoxelData, { fn decode_bencode_object(data: Object) -> Result { match data { @@ -132,13 +132,10 @@ where }; Ok(NodeContent::Internal(count as u8)) } else { - Ok(NodeContent::::Leaf(array_init::array_init(|_| { - array_init::array_init(|_| { - array_init::array_init(|_| { - NodeContent::::decode_single(&mut list).unwrap() - }) - }) - }))) + Ok(NodeContent::::Leaf( + [[[NodeContent::::decode_single(&mut list).unwrap(); DIM]; DIM]; + DIM], + )) } } Object::Bytes(b) => { @@ -244,7 +241,7 @@ where impl FromBencode for Octree where - T: PartialEq + Default + Clone + VoxelData, + T: PartialEq + Default + Clone + Copy + VoxelData, { fn decode_bencode_object(data: Object) -> Result { match data { diff --git a/src/octree/detail.rs b/src/octree/detail.rs index daeda4e..2ae2ca3 100644 --- a/src/octree/detail.rs +++ b/src/octree/detail.rs @@ -149,7 +149,7 @@ where ///#################################################################################### impl NodeContent where - T: VoxelData + PartialEq + Clone + Default, + T: VoxelData + PartialEq + Clone + Copy + Default, { pub fn is_leaf(&self) -> bool { matches!(self, NodeContent::Leaf(_)) @@ -221,9 +221,7 @@ where } pub fn leaf_from(data: T) -> Self { - NodeContent::Leaf(array_init::array_init(|_| { - array_init::array_init(|_| array_init::array_init(|_| data.clone())) - })) + NodeContent::Leaf([[[data; DIM]; DIM]; DIM]) } } diff --git a/src/octree/mod.rs b/src/octree/mod.rs index 0aa9ef7..17c9d07 100644 --- a/src/octree/mod.rs +++ b/src/octree/mod.rs @@ -18,7 +18,10 @@ use crate::octree::{ use crate::spatial::{math::hash_region, Cube}; use bendy::{decoding::FromBencode, encoding::ToBencode}; -impl Octree { +impl Octree +where + T: Default + PartialEq + Clone + Copy + VoxelData, +{ /// converts the data structure to a byte representation pub fn to_bytes(&self) -> Vec { self.to_bencode().ok().unwrap() diff --git a/src/octree/raytracing/raytracing_on_cpu.rs b/src/octree/raytracing/raytracing_on_cpu.rs index 6ae71b9..a289ee0 100644 --- a/src/octree/raytracing/raytracing_on_cpu.rs +++ b/src/octree/raytracing/raytracing_on_cpu.rs @@ -34,8 +34,9 @@ impl NodeStackItem { } } -impl - Octree +impl Octree +where + T: Default + PartialEq + Clone + Copy + VoxelData, { pub(in crate::octree) fn get_dda_scale_factors(ray: &Ray) -> V3c { V3c::new( diff --git a/src/octree/update.rs b/src/octree/update.rs index 95dab6f..29d830d 100644 --- a/src/octree/update.rs +++ b/src/octree/update.rs @@ -11,7 +11,10 @@ use crate::spatial::{ Cube, }; -impl Octree { +impl Octree +where + T: Default + PartialEq + Clone + Copy + VoxelData, +{ /// Inserts the given data into the octree into the intended voxel position pub fn insert(&mut self, position: &V3c, data: T) -> Result<(), OctreeError> { self.insert_at_lod(position, 1, data)