Resolves #45 - implementation of GPU cache

Cargo.toml

@@ -6,18 +6,19 @@ authors = ["Dávid Tóth <[email protected]>"]
 license = "MIT OR Apache-2.0"
 
 [features]
-default = ["dot_vox_support"]
+default = ["bevy_wgpu","dot_vox_support"]
 raytracing = ["dep:image", "dep:show-image"]
 serialization = ["dep:serde"]
 dot_vox_support = ["dep:dot_vox", "dep:nalgebra"]
-bevy_wgpu = ["raytracing", "dep:bevy", "dep:iyes_perf_ui"]
+bevy_wgpu = ["raytracing", "dep:bevy", "dep:iyes_perf_ui", "dep:crossbeam"]
 
 [dependencies]
 num-traits = "0.2.19"
 serde = { version = "1.0.183", features = ["derive"], optional = true }
 bendy = { git = "https://github.com/davids91/bendy.git" , features = ["std", "serde"]}
 dot_vox = { version = "5.1.1", optional = true }
 nalgebra = { version = "0.33.0", optional = true }
+crossbeam = { version = "0.8.4", optional = true }
 
 # for example cpu_render
 image = { version = "0.25.1", optional = true }

assets/shaders/viewport_render.wgsl

@@ -316,15 +316,13 @@ fn dda_step_to_next_sibling(
 }
 
 // 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::octree::raytracing::classic_raytracing_on_bevy_wgpu::set_node_meta_inner
+fn get_node_occupancy_bitmap(node_key: u32, data_meta_bytes: u32) -> u32{
+    if 0 == (node_key % 2) {
+        return ((data_meta_bytes & 0x0000FF00u) >> 8u );
+    } else {
+        return ((data_meta_bytes & 0xFF000000u) >> 24u );
+    }
 }
 
 //crate::spatial::math::step_octant
@@ -389,21 +387,21 @@ fn traverse_brick(
         - brick_bounds.min_position
     );
     current_index = vec3i(
-        clamp(i32(current_index.x), 0, i32(octreeMetaData.voxel_brick_dim - 1)),
-        clamp(i32(current_index.y), 0, i32(octreeMetaData.voxel_brick_dim - 1)),
-        clamp(i32(current_index.z), 0, i32(octreeMetaData.voxel_brick_dim - 1))
+        clamp(i32(current_index.x), 0, i32(octree_meta_data.voxel_brick_dim - 1)),
+        clamp(i32(current_index.y), 0, i32(octree_meta_data.voxel_brick_dim - 1)),
+        clamp(i32(current_index.z), 0, i32(octree_meta_data.voxel_brick_dim - 1))
     );
 
     var current_bounds = Cube(
         (
             brick_bounds.min_position 
-            + vec3f(current_index) * (brick_bounds.size / f32(octreeMetaData.voxel_brick_dim))
+            + vec3f(current_index) * (brick_bounds.size / f32(octree_meta_data.voxel_brick_dim))
         ),
-        (brick_bounds.size / f32(octreeMetaData.voxel_brick_dim))
+        (brick_bounds.size / f32(octree_meta_data.voxel_brick_dim))
     );
 
     var mapped_index = position_in_bitmap_64bits(
-        vec3u(current_index), octreeMetaData.voxel_brick_dim
+        vec3u(current_index), octree_meta_data.voxel_brick_dim
     );
     if (
         0 == (
@@ -419,21 +417,21 @@ fn traverse_brick(
     }
 
     var prev_bitmap_position_full_resolution = vec3u(
-        vec3f(current_index) * (4. / f32(octreeMetaData.voxel_brick_dim))
+        vec3f(current_index) * (4. / f32(octree_meta_data.voxel_brick_dim))
     );
     loop{
         if current_index.x < 0
-            || current_index.x >= i32(octreeMetaData.voxel_brick_dim)
+            || current_index.x >= i32(octree_meta_data.voxel_brick_dim)
             || current_index.y < 0
-            || current_index.y >= i32(octreeMetaData.voxel_brick_dim)
+            || current_index.y >= i32(octree_meta_data.voxel_brick_dim)
             || current_index.z < 0
-            || current_index.z >= i32(octreeMetaData.voxel_brick_dim)
+            || current_index.z >= i32(octree_meta_data.voxel_brick_dim)
         {
             return BrickHit(false, vec3u());
         }
 
         let bitmap_position_full_resolution = vec3u(
-            vec3f(current_index) * (4. / f32(octreeMetaData.voxel_brick_dim))
+            vec3f(current_index) * (4. / f32(octree_meta_data.voxel_brick_dim))
         );
         if(
             bitmap_position_full_resolution.x != prev_bitmap_position_full_resolution.x
@@ -458,7 +456,7 @@ fn traverse_brick(
 
         mapped_index = u32(flat_projection(
             vec3u(current_index),
-            vec2u(octreeMetaData.voxel_brick_dim, octreeMetaData.voxel_brick_dim)
+            vec2u(octree_meta_data.voxel_brick_dim, octree_meta_data.voxel_brick_dim)
         ));
         if (brick_index_start + mapped_index) >= arrayLength(&voxels)
         {
@@ -477,7 +475,7 @@ fn traverse_brick(
         );
         current_bounds.min_position = (
             current_bounds.min_position
-            + vec3f(step) * (brick_bounds.size / f32(octreeMetaData.voxel_brick_dim))
+            + vec3f(step) * (brick_bounds.size / f32(octree_meta_data.voxel_brick_dim))
         );
         current_index = current_index + vec3i(round(step));
     }
@@ -500,7 +498,7 @@ fn get_by_ray(ray: Line) -> OctreeRayIntersection{
 
     var node_stack: array<u32, NODE_STACK_SIZE>;
     var node_stack_meta: u32 = 0;
-    var current_bounds = Cube(vec3(0.), f32(octreeMetaData.octree_size));
+    var current_bounds = Cube(vec3(0.), f32(octree_meta_data.octree_size));
     var ray_current_distance: f32  = 0.0;
     var target_octant = OOB_OCTANT;
     var current_node_key = EMPTY_MARKER;
@@ -516,26 +514,34 @@ fn get_by_ray(ray: Line) -> OctreeRayIntersection{
 
     while target_octant != OOB_OCTANT {
         current_node_key = OCTREE_ROOT_NODE_KEY;
-        current_bounds = Cube(vec3(0.), f32(octreeMetaData.octree_size));
+        current_bounds = Cube(vec3(0.), f32(octree_meta_data.octree_size));
         node_stack_push(&node_stack, &node_stack_meta, OCTREE_ROOT_NODE_KEY);
         while(!node_stack_is_empty(node_stack_meta)) {
             var leaf_miss = false;
-            if is_leaf(nodes[current_node_key].sized_node_meta) {
+            if ( // is_leaf
+                (
+                    (0 == (current_node_key % 2))
+                    && (0 != (data_meta_bytes[current_node_key / 2] & 0x00000004u))
+                )||(
+                    (1 == (current_node_key % 2))
+                    && (0 != (data_meta_bytes[current_node_key / 2] & 0x00040000u))
+                )
+            ){
                 let leaf_brick_hit = traverse_brick(
                     ray, &ray_current_distance, nodes[current_node_key].voxels_start_at,
-                    children_buffer[nodes[current_node_key].children_starts_at],
-                    children_buffer[nodes[current_node_key].children_starts_at + 1],
+                    node_children[nodes[current_node_key].children_starts_at],
+                    node_children[nodes[current_node_key].children_starts_at + 1],
                     current_bounds, ray_scale_factors, direction_lut_index
                 );
                 if leaf_brick_hit.hit == true {
                     let hit_in_voxels = (
                         nodes[current_node_key].voxels_start_at
                         + u32(flat_projection(
                             leaf_brick_hit.index,
-                            vec2u(octreeMetaData.voxel_brick_dim, octreeMetaData.voxel_brick_dim)
+                            vec2u(octree_meta_data.voxel_brick_dim, octree_meta_data.voxel_brick_dim)
                         ))
                     );
-                    current_bounds.size = round(current_bounds.size / f32(octreeMetaData.voxel_brick_dim));
+                    current_bounds.size = round(current_bounds.size / f32(octree_meta_data.voxel_brick_dim));
                     current_bounds.min_position = current_bounds.min_position
                         + vec3f(leaf_brick_hit.index) * current_bounds.size;
                     return OctreeRayIntersection(
@@ -552,9 +558,9 @@ fn get_by_ray(ray: Line) -> OctreeRayIntersection{
             if( leaf_miss
                 || target_octant == OOB_OCTANT
                 || EMPTY_MARKER == current_node_key // Should never happen
-                || 0 == get_node_occupancy_bitmap(nodes[current_node_key].sized_node_meta)
+                || 0 == get_node_occupancy_bitmap(current_node_key, data_meta_bytes[current_node_key / 2])
                 || ( 0 == (
-                    get_node_occupancy_bitmap(nodes[current_node_key].sized_node_meta)
+                    get_node_occupancy_bitmap(current_node_key, data_meta_bytes[current_node_key / 2])
                     & RAY_TO_NODE_OCCUPANCY_BITMASK_LUT[target_octant][direction_lut_index]
                 ))
             ){
@@ -587,11 +593,13 @@ fn get_by_ray(ray: Line) -> OctreeRayIntersection{
             }
 
             var target_bounds = child_bounds_for(current_bounds, target_octant);
-            var target_child_key = children_buffer[nodes[current_node_key].children_starts_at + target_octant];
+            var target_child_key = node_children[
+                nodes[current_node_key].children_starts_at + target_octant
+            ];
             if (
                 target_child_key < arrayLength(&nodes) //!crate::object_pool::key_is_valid
                 && 0 != (
-                    get_node_occupancy_bitmap(nodes[current_node_key].sized_node_meta)
+                    get_node_occupancy_bitmap(current_node_key, data_meta_bytes[current_node_key / 2])
                     & ( // crate::spatial::math::octant_bitmask
                         0x00000001u << (target_octant & 0x000000FF)
                     )
@@ -617,7 +625,7 @@ fn get_by_ray(ray: Line) -> OctreeRayIntersection{
                     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[
+                        target_child_key = node_children[
                             nodes[current_node_key].children_starts_at + target_octant
                         ];
                     }
@@ -627,11 +635,11 @@ fn get_by_ray(ray: Line) -> OctreeRayIntersection{
                         || (
                             target_child_key < arrayLength(&nodes) //crate::object_pool::key_is_valid
                             && 0 != (
-                                get_node_occupancy_bitmap(nodes[current_node_key].sized_node_meta)
+                                get_node_occupancy_bitmap(current_node_key, data_meta_bytes[current_node_key / 2])
                                 & (0x00000001u << target_octant) // crate::spatial::math::octant_bitmask
                             )
                             && 0 != (
-                                get_node_occupancy_bitmap(nodes[target_child_key].sized_node_meta)
+                                get_node_occupancy_bitmap(target_child_key, data_meta_bytes[target_child_key / 2])
                                 & RAY_TO_NODE_OCCUPANCY_BITMASK_LUT[hash_region(
                                     point_in_ray_at_distance(ray, ray_current_distance) - target_bounds.min_position,
                                     round(target_bounds.size / 2.)
@@ -651,15 +659,15 @@ fn get_by_ray(ray: Line) -> OctreeRayIntersection{
             + step_vec * current_bounds.size
         );
         if(
-          current_octant_center.x < f32(octreeMetaData.octree_size)
-          && current_octant_center.y < f32(octreeMetaData.octree_size)
-          && current_octant_center.z < f32(octreeMetaData.octree_size)
+          current_octant_center.x < f32(octree_meta_data.octree_size)
+          && current_octant_center.y < f32(octree_meta_data.octree_size)
+          && current_octant_center.z < f32(octree_meta_data.octree_size)
           && current_octant_center.x > 0.
           && current_octant_center.y > 0.
           && current_octant_center.z > 0.
         ) {
             target_octant = hash_region(
-                current_octant_center, f32(octreeMetaData.octree_size) / 2.
+                current_octant_center, f32(octree_meta_data.octree_size) / 2.
             );
         } else {
             target_octant = OOB_OCTANT;
@@ -684,7 +692,6 @@ fn is_empty(e: Voxelement) -> bool {
 }
 
 struct SizedNode {
-    sized_node_meta: u32,
     children_starts_at: u32,
     voxels_start_at: u32,
 }
@@ -710,20 +717,23 @@ var output_texture: texture_storage_2d<rgba8unorm, read_write>;
 var<uniform> viewport: Viewport;
 
 @group(1) @binding(0)
-var<uniform> octreeMetaData: OctreeMetaData;
+var<uniform> octree_meta_data: OctreeMetaData;
 
 @group(1) @binding(1)
 var<storage, read_write> nodes: array<SizedNode>;
 
 @group(1) @binding(2)
-var<storage, read_write> children_buffer: array<u32>;
+var<storage, read_write> node_children: array<u32>;
 
 @group(1) @binding(3)
 var<storage, read_write> voxels: array<Voxelement>;
 
 @group(1) @binding(4)
 var<storage, read_write> color_palette: array<vec4f>;
 
+@group(1) @binding(5)
+var<storage, read_write> data_meta_bytes: array<u32>;
+
 @compute @workgroup_size(8, 8, 1)
 fn update(
     @builtin(global_invocation_id) invocation_id: vec3<u32>,
@@ -762,12 +772,12 @@ fn update(
     /*// +++ DEBUG +++
     // Display the xyz axes
     let root_hit = cube_intersect_ray(
-        Cube(vec3(0.,0.,0.), f32(octreeMetaData.octree_size)), ray
+        Cube(vec3(0.,0.,0.), f32(octree_meta_data.octree_size)), ray
     );
     if root_hit.hit == true {
         if root_hit. impact_hit == true {
-            let axes_length = f32(octreeMetaData.octree_size) / 2.;
-            let axes_width = f32(octreeMetaData.octree_size) / 50.;
+            let axes_length = f32(octree_meta_data.octree_size) / 2.;
+            let axes_width = f32(octree_meta_data.octree_size) / 50.;
             let entry_point = point_in_ray_at_distance(ray, root_hit.impact_distance);
             if entry_point.x < axes_length && entry_point.y < axes_width && entry_point.z < axes_width {
                 rgb_result.r = 1.;

examples/minecraft.rs

@@ -7,7 +7,8 @@ use bevy::{prelude::*, window::WindowPlugin};
 #[cfg(feature = "bevy_wgpu")]
 use shocovox_rs::octree::{
     raytracing::{
-        bevy::create_viewing_glass, ShocoVoxRenderPlugin, ShocoVoxViewingGlass, Viewport,
+        bevy::create_viewing_glass, ShocoVoxRenderData, ShocoVoxRenderPlugin, ShocoVoxViewingGlass,
+        Viewport,
     },
     Albedo, V3c,
 };
@@ -64,11 +65,6 @@ fn setup(mut commands: Commands, images: ResMut<Assets<Image>>) {
         tree.save(&tree_path).ok().unwrap();
     }
 
-    let origin = V3c::new(
-        tree.get_size() as f32 * 2.,
-        tree.get_size() as f32 / 2.,
-        tree.get_size() as f32 * -2.,
-    );
     commands.spawn(DomePosition {
         yaw: 0.,
         roll: 0.,
@@ -152,9 +148,14 @@ fn rotate_camera(
 #[cfg(feature = "bevy_wgpu")]
 fn handle_zoom(
     keys: Res<ButtonInput<KeyCode>>,
-    mut viewing_glass: ResMut<ShocoVoxViewingGlass>,
     mut angles_query: Query<&mut DomePosition>,
+    svx_data: Option<ResMut<ShocoVoxRenderData>>,
 ) {
+    if let Some(ref svx_data) = svx_data {
+        if svx_data.is_changed() {
+            println!("changed!! --> {:?}", svx_data.node_children[0..2].to_vec());
+        }
+    }
     const ADDITION: f32 = 0.05;
     let angle_update_fn = |angle, delta| -> f32 {
         let new_angle = angle + delta;
@@ -184,6 +185,9 @@ fn handle_zoom(
     if keys.pressed(KeyCode::PageDown) {
         angles_query.single_mut().radius *= 1.1;
     }
+    if keys.pressed(KeyCode::Tab) {
+        svx_data.unwrap().do_the_thing = true;
+    }
 }
 
 #[cfg(not(feature = "bevy_wgpu"))]

src/octree/mod.rs

@@ -81,7 +81,7 @@ where
     /// Provides immutable reference to the data, if there is any at the given position
     pub fn get(&self, position: &V3c<u32>) -> Option<&T> {
         let mut current_bounds = Cube::root_bounds(self.octree_size as f32);
-        let mut current_node_key = Octree::<T, DIM>::ROOT_NODE_KEY as usize;
+        let mut current_node_key = Self::ROOT_NODE_KEY as usize;
         let position = V3c::from(*position);
         if !bound_contains(&current_bounds, &position) {
             return None;
@@ -118,7 +118,7 @@ where
     /// Provides mutable reference to the data, if there is any at the given position
     pub fn get_mut(&mut self, position: &V3c<u32>) -> Option<&mut T> {
         let mut current_bounds = Cube::root_bounds(self.octree_size as f32);
-        let mut current_node_key = Octree::<T, DIM>::ROOT_NODE_KEY as usize;
+        let mut current_node_key = Self::ROOT_NODE_KEY as usize;
         let position = V3c::from(*position);
         if !bound_contains(&current_bounds, &position) {
             return None;