diff --git a/Cargo.toml b/Cargo.toml
index 4dd0821..a2d7df9 100644
--- a/Cargo.toml
+++ b/Cargo.toml
@@ -1,16 +1,16 @@
 [package]
 name = "shocovox-rs"
-version = "0.4.1"
+version = "0.5.0"
 edition = "2021"
 authors = ["Dávid Tóth <toth.david.munka@gmail.com>"]
 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", "dep:bimap"]
 
 [dependencies]
 num-traits = "0.2.19"
@@ -18,6 +18,8 @@ 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 }
+bimap = { version = "0.6.3", optional = true }
 
 # for example cpu_render
 image = { version = "0.25.1", optional = true }
diff --git a/README.md b/README.md
index 78374ed..088ed7f 100644
--- a/README.md
+++ b/README.md
@@ -12,13 +12,18 @@ Roadmap:
 
 Issue spotlight: 
 -
-These are the issues I will work on until 2025 Q1. I am eliminating them in a sequential manner.
-- #54 - Occupied bitmap structure update - Restructure of the per-node and per-brick occupied bitmaps to be more efficient(gotta go fast)
-- #56 - Rework user data handling - Trimming the fat in Voxel storage, broadening the possibilities with user data
-- #45 - GPU cache - To make it possible to display octrees of limitless size on the GPU by streaming only what one can see
+These are the issues I will work on until 2025 Q2. I am eliminating them in a sequential manner.
+- #56 - Introduce Palettes - Trimming the fat in Voxel storage, broadening the possibilities with user data and eliminating some data conversion overhead with bricks.
+- #65 - Flatten brick storage: trimming some additional overhead, and eliminating some possible techDebt (`DIM` generic argument with Octree)
 - #3 - to make it possible to have a limitless octree: so it's not bound by the RAM size anymore
+- #17 Beam Optimization - Pre-render a small resolution image to optimally initialize ray distances, and help with deciding which bricks to load pre-emptively. GOTTA GO FAST
 - #28, #6 - Level of Detail implementation to render large scenes more efficiently
 
+If you feel adventurous:
+-
+
+I have marked some issues with the help needed flag, which I think would be a good addition to the library, but I can not focus on it as I am but a single person with limited time and resources. Feel free to try to tackle any of the marked issues (Or basically anything you'd like), I will provide any needed help and support if I can. 
+
 Special thanks to contributors and supporters!
 -
 
diff --git a/assets/shaders/viewport_render.wgsl b/assets/shaders/viewport_render.wgsl
index 02f2c57..96f714d 100644
--- a/assets/shaders/viewport_render.wgsl
+++ b/assets/shaders/viewport_render.wgsl
@@ -234,6 +234,73 @@ fn dda_step_to_next_sibling(
     return result;
 }
 
+// Unique to this implementation, not adapted from rust code
+/// Sets the used bit true for the given node
+fn set_node_used(node_key: u32) {
+    if 0 != (metadata[node_key] & 0x01u) {
+        // no need to set if already true
+        return;
+    }
+
+    loop{
+        let exchange_result = atomicCompareExchangeWeak(
+            &metadata[node_key], metadata[node_key], metadata[node_key] | 0x01u
+        );
+        if(exchange_result.exchanged || 0 < (exchange_result.old_value & 0x01u)){
+            break;
+        }
+    }
+}
+
+// Unique to this implementation, not adapted from rust code
+/// Sets the used bit true for the given brick
+fn set_brick_used(brick_index: u32) {
+    if 0 != ( metadata[brick_index / 8] & (0x01u << (24u + (brick_index % 8u))) ) {
+        // no need to set if already true
+        return;
+    }
+
+    loop{
+        let exchange_result = atomicCompareExchangeWeak(
+            &metadata[brick_index / 8],
+            metadata[brick_index / 8],
+            metadata[brick_index / 8] | (0x01u << (24u + (brick_index % 8u)))
+        );
+        if(
+            exchange_result.exchanged
+            || 0 != ( exchange_result.old_value & (0x01u << (24u + (brick_index % 8u))) )
+        ){
+            break;
+        }
+    }
+}
+
+// Unique to this implementation, not adapted from rust code
+/// Requests the child of the given node to be uploaded
+fn request_node(node_meta_index: u32, child_octant: u32) -> bool {
+    var request_index = 0u;
+    loop{
+        let exchange_result = atomicCompareExchangeWeak(
+            &node_requests[request_index], EMPTY_MARKER,
+            (node_meta_index & 0x00FFFFFFu)|((child_octant & 0x000000FF) << 24)
+        );
+        if(
+            exchange_result.exchanged 
+            ||(
+                exchange_result.old_value
+                == ((node_meta_index & 0x00FFFFFFu)|((child_octant & 0x000000FF) << 24))
+            )
+        ) {
+            break;
+        }
+        request_index += 1u;
+        if(request_index >= arrayLength(&node_requests)) {
+            return false;
+        }
+    }
+    return true;
+}
+
 //crate::spatial::math::step_octant
 fn step_octant(octant: u32, step: ptr<function, vec3f>) -> u32 {
     return (
@@ -270,7 +337,7 @@ fn traverse_brick(
     ray_scale_factors: ptr<function, vec3f>,
     direction_lut_index: u32,
 ) -> BrickHit {
-    let dimension = i32(octreeMetaData.voxel_brick_dim);
+    let dimension = i32(octree_meta_data.voxel_brick_dim);
 
     var current_index = clamp(
         vec3i(vec3f( // entry position in brick
@@ -352,13 +419,14 @@ fn probe_brick(
     ray_scale_factors: ptr<function, vec3f>,
     direction_lut_index: u32,
 ) -> OctreeRayIntersection {
-    if(0 != ((0x01u << (8 + brick_octant)) & nodes[leaf_node_key])) { // brick is not empty
-        let brick_start_index = children_buffer[((leaf_node_key * 8) + brick_octant)];
-        if(0 == ((0x01u << (16 + brick_octant)) & nodes[leaf_node_key])) { // brick is solid
+    if(0 != ((0x01u << (8 + brick_octant)) & metadata[leaf_node_key])) { // brick is not empty
+        let brick_index = node_children[((leaf_node_key * 8) + brick_octant)];
+        set_brick_used(brick_index);
+        if(0 == ((0x01u << (16 + brick_octant)) & metadata[leaf_node_key])) { // brick is solid
             // Whole brick is solid, ray hits it at first connection
             return OctreeRayIntersection(
                 true,
-                color_palette[brick_start_index], // Albedo is in color_palette
+                color_palette[brick_index], // Albedo is in color_palette, data is not a brick index in this case
                 0, // user data lost for now as color palette doesn't have it.. sorry
                 point_in_ray_at_distance(ray, *ray_current_distance),
                 cube_impact_normal(*brick_bounds, point_in_ray_at_distance(ray, *ray_current_distance))
@@ -366,7 +434,7 @@ fn probe_brick(
         } else { // brick is parted
             let leaf_brick_hit = traverse_brick(
                 ray, ray_current_distance,
-                brick_start_index,
+                brick_index,
                 brick_bounds, ray_scale_factors, direction_lut_index
             );
             if leaf_brick_hit.hit == true {
@@ -379,9 +447,9 @@ fn probe_brick(
                         Cube(
                             (*brick_bounds).min_position + (
                                 vec3f(leaf_brick_hit.index)
-                                * round((*brick_bounds).size / f32(octreeMetaData.voxel_brick_dim))
+                                * round((*brick_bounds).size / f32(octree_meta_data.voxel_brick_dim))
                             ),
-                            round((*brick_bounds).size / f32(octreeMetaData.voxel_brick_dim)),
+                            round((*brick_bounds).size / f32(octree_meta_data.voxel_brick_dim)),
                         ),
                         point_in_ray_at_distance(ray, *ray_current_distance)
                     )
@@ -393,18 +461,93 @@ fn probe_brick(
     return OctreeRayIntersection(false, vec4f(0.), 0, vec3f(0.), vec3f(0., 0., 1.));
 }
 
-fn get_by_ray(ray: ptr<function, Line>) -> OctreeRayIntersection{
+// Unique to this implementation, not adapted from rust code
+/// Traverses the node to provide information about how the occupied bits of the node
+/// and the given ray collides. The higher the number, the closer the hit is.
+fn traverse_node_for_ocbits(
+    ray: ptr<function, Line>,
+    ray_current_distance: ptr<function,f32>,
+    node_key: u32,
+    node_bounds: ptr<function, Cube>,
+    ray_scale_factors: ptr<function, vec3f>,
+) -> f32 {
+    let original_distance = *ray_current_distance;
+
+    var position = vec3f(
+        point_in_ray_at_distance(ray, *ray_current_distance)
+        - (*node_bounds).min_position
+    );
+
+    var current_index = vec3i(vec3f(
+        clamp( (position.x * 4. / (*node_bounds).size), 0.5, 3.5),
+        clamp( (position.y * 4. / (*node_bounds).size), 0.5, 3.5),
+        clamp( (position.z * 4. / (*node_bounds).size), 0.5, 3.5),
+    ));
+
+    var current_bounds = Cube(
+        (
+            (*node_bounds).min_position
+            + vec3f(current_index) * ((*node_bounds).size / 4.)
+        ),
+        round((*node_bounds).size / 4.)
+    );
+
+    var safety = 0u;
+    var result = 0.;
+    loop {
+        if safety > 10 || current_index.x < 0 || current_index.x >= 4
+            || current_index.y < 0 || current_index.y >= 4
+            || current_index.z < 0 || current_index.z >= 4
+        {
+            break;
+        }
+
+        let bitmap_index = BITMAP_INDEX_LUT[u32(current_index.x)]
+                                           [u32(current_index.y)]
+                                           [u32(current_index.z)];
+        if (
+            (
+                (bitmap_index < 32)
+                && (0u != (node_occupied_bits[node_key * 2]
+                            & (0x01u << bitmap_index) ))
+            )||(
+                (bitmap_index >= 32)
+                && (0u != (node_occupied_bits[node_key * 2 + 1]
+                            & (0x01u << (bitmap_index - 32)) ))
+            )
+        ){
+            result = 1. - (f32(safety) * 0.25);
+            break;
+        }
+
+        let step = round(dda_step_to_next_sibling(
+            ray,
+            ray_current_distance,
+            &current_bounds,
+            ray_scale_factors
+        ));
+        current_bounds.min_position += step * current_bounds.size;
+        current_index += vec3i(step);
+        safety += 1u;
+    }
+
+    *ray_current_distance = original_distance;
+    return result;
+}
+
+fn get_by_ray(ray: ptr<function, Line>) -> OctreeRayIntersection {
     var ray_scale_factors = get_dda_scale_factors(ray); // Should be const, but then it can't be passed as ptr
     let direction_lut_index = hash_direction((*ray).direction);
 
     var node_stack: array<u32, NODE_STACK_SIZE>;
     var node_stack_meta: u32 = 0;
     var ray_current_distance = 0.0;
-    var current_bounds = Cube(vec3(0.), f32(octreeMetaData.octree_size));
+    var current_bounds = Cube(vec3(0.), f32(octree_meta_data.octree_size));
     var current_node_key = EMPTY_MARKER;
     var current_node_meta = 0u;
     var target_octant = OOB_OCTANT;
     var step_vec = vec3f(0.);
+    var missing_data_color = vec3f(0.);
 
     let root_intersect = cube_intersect_ray(current_bounds, ray);
     if(root_intersect.hit){
@@ -422,15 +565,15 @@ fn get_by_ray(ray: ptr<function, Line>) -> OctreeRayIntersection{
     while target_octant != OOB_OCTANT {
         /*// +++ DEBUG +++
         outer_safety += 1;
-        if(f32(outer_safety) > f32(octreeMetaData.octree_size) * sqrt(3.)) {
+        if(f32(outer_safety) > f32(octree_meta_data.octree_size) * sqrt(3.)) {
             return OctreeRayIntersection(
                 true, vec4f(1.,0.,0.,1.), 0, vec3f(0.), vec3f(0., 0., 1.)
             );
         }
        */ // --- DEBUG ---
         current_node_key = OCTREE_ROOT_NODE_KEY;
-        current_node_meta = nodes[OCTREE_ROOT_NODE_KEY];
-        current_bounds = Cube(vec3(0.), f32(octreeMetaData.octree_size));
+        current_node_meta = metadata[OCTREE_ROOT_NODE_KEY];
+        current_bounds = Cube(vec3(0.), f32(octree_meta_data.octree_size));
         node_stack_push(&node_stack, &node_stack_meta, OCTREE_ROOT_NODE_KEY);
         /*// +++ DEBUG +++
         var safety = 0;
@@ -438,50 +581,107 @@ fn get_by_ray(ray: ptr<function, Line>) -> OctreeRayIntersection{
         while(!node_stack_is_empty(node_stack_meta)) {
             /*// +++ DEBUG +++
             safety += 1;
-            if(f32(safety) > f32(octreeMetaData.octree_size) * sqrt(30.)) {
+            if(f32(safety) > f32(octree_meta_data.octree_size) * sqrt(30.)) {
                 return OctreeRayIntersection(
                     true, vec4f(0.,0.,1.,1.), 0, vec3f(0.), vec3f(0., 0., 1.)
                 );
             }
             */// --- DEBUG ---
-            var target_bounds: Cube;
             var do_backtrack_after_leaf_miss = false;
+            var target_child_key = node_children[(current_node_key * 8) + target_octant];
+            var target_bounds = child_bounds_for(&current_bounds, target_octant);
+            var bitmap_pos_in_node = clamp(
+                (
+                    point_in_ray_at_distance(ray, ray_current_distance)
+                    - current_bounds.min_position
+                ) * 4. / current_bounds.size,
+                vec3f(FLOAT_ERROR_TOLERANCE),
+                vec3f(4. - FLOAT_ERROR_TOLERANCE)
+            );
+
+            if(
+                // In case node doesn't yet have the target child node uploaded to GPU
+                (0 == (0x00000004 & current_node_meta)) // node is not a leaf
+                && target_octant != OOB_OCTANT
+                && target_child_key == EMPTY_MARKER // target child key is invalid
+                && ( // node is occupied at target octant
+                    0 != (
+                        BITMAP_MASK_FOR_OCTANT_LUT[target_octant][0]
+                        & node_occupied_bits[current_node_key * 2]
+                    )
+                    || 0 != (
+                        BITMAP_MASK_FOR_OCTANT_LUT[target_octant][1]
+                        & node_occupied_bits[current_node_key * 2 + 1]
+                    )
+                )
+                // Request node only once per ray iteration to prioritize nodes in sight for cache
+                && 0 == (missing_data_color.r + missing_data_color.g + missing_data_color.b)
+            ){
+                if request_node(current_node_key, target_octant) {
+                    missing_data_color += (
+                        vec3f(0.5,0.3,0.0) *
+                        vec3f(traverse_node_for_ocbits(
+                            ray,
+                            &ray_current_distance,
+                            current_node_key,
+                            &current_bounds,
+                            &ray_scale_factors
+                        ))
+                    );
+                } else {
+                    missing_data_color += (
+                        vec3f(0.7,0.2,0.0) *
+                        vec3f(traverse_node_for_ocbits(
+                            ray,
+                            &ray_current_distance,
+                            current_node_key,
+                            &current_bounds,
+                            &ray_scale_factors
+                        ))
+                    );
+                }
+            }
+
             if (target_octant != OOB_OCTANT) {
                 if(0 != (0x00000004 & current_node_meta)) { // node is leaf
                     var hit: OctreeRayIntersection;
-                    if(0 != (0x00000008 & current_node_meta)) { // node is a uniform leaf
-                        hit = probe_brick(
-                            ray, &ray_current_distance,
-                            current_node_key, 0u, &current_bounds,
-                            &ray_scale_factors, direction_lut_index
-                        );
-                        if hit.hit == true {
-                            return hit;
+
+                    if // node not empty at target octant, while the brick is marked unavailable
+                        (0 != ((0x01u << (8 + target_octant)) & current_node_meta))
+                        && EMPTY_MARKER == node_children[(current_node_key * 8) + target_octant]
+                    {
+                        // child brick is not yet uploaded to GPU
+                        if request_node(current_node_key, target_octant) {
+                            missing_data_color += vec3f(0.3,0.1,0.0);
+                        } else {
+                            missing_data_color += vec3f(0.7,0.0,0.0);
                         }
                         do_backtrack_after_leaf_miss = true;
-                    } else { // node is a non-uniform leaf
-                        target_bounds = child_bounds_for(&current_bounds, target_octant);
-                        hit = probe_brick(
-                            ray, &ray_current_distance,
-                            current_node_key, target_octant,
-                            &target_bounds,
-                            &ray_scale_factors, direction_lut_index
-                        );
+                    } else {
+                        if(0 != (0x00000008 & current_node_meta)) { // node is a uniform leaf
+                            hit = probe_brick(
+                                ray, &ray_current_distance,
+                                current_node_key, 0u, &current_bounds,
+                                &ray_scale_factors, direction_lut_index
+                            );
+                            do_backtrack_after_leaf_miss = true;
+                        } else { // node is a non-uniform leaf
+                            target_bounds = child_bounds_for(&current_bounds, target_octant);
+                            hit = probe_brick(
+                                ray, &ray_current_distance,
+                                current_node_key, target_octant,
+                                &target_bounds,
+                                &ray_scale_factors, direction_lut_index
+                            );
+                        }
                         if hit.hit == true {
+                            hit.albedo += vec4f(missing_data_color, 0.);
                             return hit;
                         }
                     }
                 }
             }
 
-            var bitmap_pos_in_node = clamp(
-                (
-                    point_in_ray_at_distance(ray, ray_current_distance)
-                    - current_bounds.min_position
-                ) * 4. / current_bounds.size,
-                vec3f(FLOAT_ERROR_TOLERANCE),
-                vec3f(4. - FLOAT_ERROR_TOLERANCE)
-            );
             if( do_backtrack_after_leaf_miss
                 || target_octant == OOB_OCTANT
                 || EMPTY_MARKER == current_node_key // Guards statements in other conditions, but should never happen
@@ -503,7 +703,7 @@ fn get_by_ray(ray: ptr<function, Line>) -> OctreeRayIntersection{
                         & node_occupied_bits[current_node_key * 2 + 1]
                     )
                 )
-            ){
+            ) {
                 // POP
                 node_stack_pop(&node_stack, &node_stack_meta);
                 step_vec = dda_step_to_next_sibling(
@@ -513,7 +713,7 @@ fn get_by_ray(ray: ptr<function, Line>) -> OctreeRayIntersection{
                 );
                 if(EMPTY_MARKER != node_stack_last(node_stack_meta)){
                     current_node_key = node_stack[node_stack_last(node_stack_meta)];
-                    current_node_meta = nodes[current_node_key];
+                    current_node_meta = metadata[current_node_key];
                     target_octant = step_octant(
                         hash_region( // parent current target octant
                             // current bound center
@@ -532,12 +732,10 @@ fn get_by_ray(ray: ptr<function, Line>) -> OctreeRayIntersection{
                 continue;
             }
 
-            target_bounds = child_bounds_for(&current_bounds, target_octant);
-            var target_child_key = children_buffer[(current_node_key * 8) + target_octant];
             if (
                 (
                     0 == (0x00000004 & current_node_meta) // node is not a leaf
-                    && target_child_key < arrayLength(&nodes) //!crate::object_pool::key_is_valid
+                    && target_child_key != EMPTY_MARKER //crate::object_pool::key_is_valid
                 )
                 && ( // There is overlap in node occupancy and potential ray hit area
                     0 != (
@@ -559,8 +757,9 @@ fn get_by_ray(ray: ptr<function, Line>) -> OctreeRayIntersection{
                 )
             ) {
                 // PUSH
+                set_node_used(target_child_key);
                 current_node_key = target_child_key;
-                current_node_meta = nodes[current_node_key];
+                current_node_meta = metadata[current_node_key];
                 current_bounds = target_bounds;
                 target_octant = hash_region( // child_target_octant
                     (point_in_ray_at_distance(ray, ray_current_distance) - target_bounds.min_position),
@@ -589,13 +788,13 @@ fn get_by_ray(ray: ptr<function, 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[(current_node_key * 8) + target_octant];
+                        target_child_key = node_children[(current_node_key * 8) + target_octant];
                         bitmap_pos_in_node += step_vec * 4. / current_bounds.size;
                     }
                     if (
                         target_octant == OOB_OCTANT
                         || ( // In case the current internal node has a valid target child
-                            target_child_key < arrayLength(&nodes) //crate::object_pool::key_is_valid
+                            target_child_key != EMPTY_MARKER //crate::object_pool::key_is_valid
                             && 0 == (0x00000004 & current_node_meta) // node is not a leaf
                             && ( // target child is in the area the ray can potentially hit
                                 0 != (
@@ -633,21 +832,21 @@ fn get_by_ray(ray: ptr<function, 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;
         }
     } // while (ray inside root bounds)
-    return OctreeRayIntersection(false, vec4f(0.), 0, vec3f(0.), vec3f(0., 0., 1.));
+    return OctreeRayIntersection(false, vec4f(missing_data_color, 1.), 0, vec3f(0.), vec3f(0., 0., 1.));
 }
 
 struct Voxelement {
@@ -685,24 +884,28 @@ var output_texture: texture_storage_2d<rgba8unorm, read_write>;
 @group(0) @binding(1)
 var<uniform> viewport: Viewport;
 
+@group(0) @binding(2)
+var<storage, read_write> node_requests: array<atomic<u32>>;
+
 @group(1) @binding(0)
-var<uniform> octreeMetaData: OctreeMetaData;
+var<uniform> octree_meta_data: OctreeMetaData;
 
 @group(1) @binding(1)
-var<storage, read_write> nodes: array<u32>;
+var<storage, read_write> metadata: array<atomic<u32>>;
 
 @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>;
+var<storage, read_write> node_occupied_bits: array<u32>;
 
 @group(1) @binding(4)
-var<storage, read_write> node_occupied_bits: array<u32>;
+var<storage, read_write> voxels: array<Voxelement>;
 
 @group(1) @binding(5)
 var<storage, read_write> color_palette: array<vec4f>;
 
+
 @compute @workgroup_size(8, 8, 1)
 fn update(
     @builtin(global_invocation_id) invocation_id: vec3<u32>,
@@ -737,16 +940,18 @@ fn update(
                 dot(ray_result.impact_normal, vec3f(-0.5,0.5,-0.5)) / 2. + 0.5
             )
         ).rgb;
+    } else {
+        rgb_result = (rgb_result + ray_result.albedo.rgb) / 2.;
     }
 
     /*// +++ DEBUG +++
-    var root_bounds = Cube(vec3(0.,0.,0.), f32(octreeMetaData.octree_size));
+    var root_bounds = Cube(vec3(0.,0.,0.), f32(octree_meta_data.octree_size));
     let root_intersect = cube_intersect_ray(root_bounds, &ray);
     if root_intersect.hit == true {
         // Display the xyz axes
         if root_intersect. 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_intersect.impact_distance);
             if entry_point.x < axes_length && entry_point.y < axes_width && entry_point.z < axes_width {
                 rgb_result.r = 1.;
@@ -758,10 +963,9 @@ fn update(
                 rgb_result.b = 1.;
             }
         }
-        //rgb_result.b += 0.1; // Also color in the area of the octree
+        rgb_result.b += 0.1; // Also color in the area of the octree
     }
     */// --- DEBUG ---
-
     textureStore(output_texture, vec2u(invocation_id.xy), vec4f(rgb_result, 1.));
 }
 
@@ -817,6 +1021,18 @@ var<private> BITMAP_INDEX_LUT: array<array<array<u32, 4>, 4>, 4> = array<array<a
     ),
 );
 
+//crate::spatial::math::mask_for_octant_64_bits
+var<private> BITMAP_MASK_FOR_OCTANT_LUT: array<array<u32, 2>, 8> = array<array<u32, 2>, 8>(
+    array<u32, 2>(0x00330033u,0x00000000u),
+    array<u32, 2>(0x00CC00CCu,0x00000000u),
+    array<u32, 2>(0x00000000u,0x00330033u),
+    array<u32, 2>(0x00000000u,0x00CC00CCu),
+    array<u32, 2>(0x33003300u,0x00000000u),
+    array<u32, 2>(0xCC00CC00u,0x00000000u),
+    array<u32, 2>(0x00000000u,0x33003300u),
+    array<u32, 2>(0x00000000u,0xCC00CC00u),
+);
+
 // Note: should be const
 var<private> RAY_TO_NODE_OCCUPANCY_BITMASK_LUT: array<array<u32, 16>, 64> = array<array<u32, 16>, 64>(
     array<u32, 16>(1,0,15,0,65537,65537,983055,983055,4369,0,65535,0,286331153,286331153,4294967295,4294967295,),
diff --git a/examples/dot_cube.rs b/examples/dot_cube.rs
index b4d02ad..4516c7d 100644
--- a/examples/dot_cube.rs
+++ b/examples/dot_cube.rs
@@ -9,9 +9,7 @@ use iyes_perf_ui::{
 
 #[cfg(feature = "bevy_wgpu")]
 use shocovox_rs::octree::{
-    raytracing::{
-        bevy::create_viewing_glass, ShocoVoxRenderPlugin, ShocoVoxViewingGlass, Viewport,
-    },
+    raytracing::{OctreeGPUHost, OctreeGPUView, Ray, SvxViewSet, Viewport},
     Albedo, V3c,
 };
 
@@ -30,9 +28,9 @@ fn main() {
         .insert_resource(ClearColor(Color::BLACK))
         .add_plugins((
             DefaultPlugins.set(WindowPlugin::default()),
-            ShocoVoxRenderPlugin {
-                resolution: DISPLAY_RESOLUTION,
-            },
+            shocovox_rs::octree::raytracing::RenderBevyPlugin::<Albedo, BRICK_DIMENSION>::new(
+                DISPLAY_RESOLUTION,
+            ),
             bevy::diagnostic::FrameTimeDiagnosticsPlugin,
             PerfUiPlugin,
         ))
@@ -44,8 +42,6 @@ fn main() {
 
 #[cfg(feature = "bevy_wgpu")]
 fn setup(mut commands: Commands, images: ResMut<Assets<Image>>) {
-    // use shocovox_rs::octree::raytracing::bevy::create_viewing_glass;
-
     let origin = V3c::new(
         TREE_SIZE as f32 * 2.,
         TREE_SIZE as f32 / 2.,
@@ -57,19 +53,6 @@ fn setup(mut commands: Commands, images: ResMut<Assets<Image>>) {
         .ok()
         .unwrap();
 
-    // +++ DEBUG +++
-    // tree.insert(&V3c::new(0, 0, 0), Albedo::from(0x66FFFF))
-    //     .ok()
-    //     .unwrap();
-    // tree.insert(&V3c::new(3, 3, 3), Albedo::from(0x66FFFF))
-    //     .ok()
-    //     .unwrap();
-    // assert!(tree.get(&V3c::new(3, 3, 3)).is_some());
-    // tree.insert_at_lod(&V3c::new(0, 0, 0), 128, Albedo::from(0x66FFFF))
-    //     .ok()
-    //     .unwrap();
-
-    // ---DEBUG ---
     for x in 0..TREE_SIZE {
         for y in 0..TREE_SIZE {
             for z in 0..TREE_SIZE {
@@ -109,9 +92,18 @@ fn setup(mut commands: Commands, images: ResMut<Assets<Image>>) {
         }
     }
 
-    let render_data = tree.create_bevy_view();
-    let viewing_glass = create_viewing_glass(
-        &Viewport {
+    commands.spawn(DomePosition {
+        yaw: 0.,
+        roll: 0.,
+        radius: tree.get_size() as f32 * 2.2,
+    });
+
+    let mut host = OctreeGPUHost { tree };
+    let mut views = SvxViewSet::default();
+    let output_texture = host.create_new_view(
+        &mut views,
+        40,
+        Viewport {
             origin,
             direction: (V3c::new(0., 0., 0.) - origin).normalized(),
             w_h_fov: V3c::new(10., 10., 3.),
@@ -119,24 +111,17 @@ fn setup(mut commands: Commands, images: ResMut<Assets<Image>>) {
         DISPLAY_RESOLUTION,
         images,
     );
-
-    commands.spawn(DomePosition {
-        yaw: 0.,
-        roll: 0.,
-        radius: tree.get_size() as f32 * 2.2,
-    });
+    commands.insert_resource(host);
+    commands.insert_resource(views);
     commands.spawn(SpriteBundle {
         sprite: Sprite {
             custom_size: Some(Vec2::new(1024., 768.)),
             ..default()
         },
-        texture: viewing_glass.output_texture.clone(),
+        texture: output_texture,
         ..default()
     });
     commands.spawn(Camera2dBundle::default());
-    commands.insert_resource(render_data);
-    commands.insert_resource(viewing_glass);
-
     commands.spawn((
         PerfUiRoot::default(),
         PerfUiEntryFPS {
@@ -163,31 +148,28 @@ struct DomePosition {
     yaw: f32,
     roll: f32,
 }
-
 #[cfg(feature = "bevy_wgpu")]
-fn rotate_camera(
-    angles_query: Query<&mut DomePosition>,
-    mut viewing_glass: ResMut<ShocoVoxViewingGlass>,
-) {
+fn rotate_camera(angles_query: Query<&mut DomePosition>, mut tree_view: ResMut<OctreeGPUView>) {
     let (yaw, roll) = (angles_query.single().yaw, angles_query.single().roll);
-
     let radius = angles_query.single().radius;
-    viewing_glass.viewport.origin = V3c::new(
+
+    tree_view.spyglass.viewport.origin = V3c::new(
         radius / 2. + yaw.sin() * radius,
         radius + roll.sin() * radius * 2.,
         radius / 2. + yaw.cos() * radius,
     );
-    viewing_glass.viewport.direction =
-        (V3c::unit(radius / 2.) - viewing_glass.viewport.origin).normalized();
+    tree_view.spyglass.viewport.direction =
+        (V3c::unit(radius / 2.) - tree_view.spyglass.viewport.origin).normalized();
 }
 
 #[cfg(feature = "bevy_wgpu")]
 fn handle_zoom(
     keys: Res<ButtonInput<KeyCode>>,
-    mut viewing_glass: ResMut<ShocoVoxViewingGlass>,
+    mut tree_view: ResMut<OctreeGPUView>,
     mut angles_query: Query<&mut DomePosition>,
+    tree: Res<OctreeGPUHost<Albedo, BRICK_DIMENSION>>,
 ) {
-    const ADDITION: f32 = 0.02;
+    const ADDITION: f32 = 0.05;
     let angle_update_fn = |angle, delta| -> f32 {
         let new_angle = angle + delta;
         if new_angle < 360. {
@@ -196,11 +178,74 @@ fn handle_zoom(
             0.
         }
     };
+    if keys.pressed(KeyCode::Tab) {
+        // Render the current view with CPU
+        let viewport_up_direction = V3c::new(0., 1., 0.);
+        let viewport_right_direction = viewport_up_direction
+            .cross(tree_view.spyglass.viewport.direction)
+            .normalized();
+        let pixel_width =
+            tree_view.spyglass.viewport.w_h_fov.x as f32 / DISPLAY_RESOLUTION[0] as f32;
+        let pixel_height =
+            tree_view.spyglass.viewport.w_h_fov.y as f32 / DISPLAY_RESOLUTION[1] as f32;
+        let viewport_bottom_left = tree_view.spyglass.viewport.origin
+            + (tree_view.spyglass.viewport.direction * tree_view.spyglass.viewport.w_h_fov.z)
+            - (viewport_up_direction * (tree_view.spyglass.viewport.w_h_fov.y / 2.))
+            - (viewport_right_direction * (tree_view.spyglass.viewport.w_h_fov.x / 2.));
+
+        // define light
+        let diffuse_light_normal = V3c::new(0., -1., 1.).normalized();
+
+        use image::ImageBuffer;
+        use image::Rgb;
+        let mut img = ImageBuffer::new(DISPLAY_RESOLUTION[0], DISPLAY_RESOLUTION[1]);
+
+        // cast each ray for a hit
+        for x in 0..DISPLAY_RESOLUTION[0] {
+            for y in 0..DISPLAY_RESOLUTION[1] {
+                let actual_y_in_image = DISPLAY_RESOLUTION[1] - y - 1;
+                //from the origin of the camera to the current point of the viewport
+                let glass_point = viewport_bottom_left
+                    + viewport_right_direction * x as f32 * pixel_width
+                    + viewport_up_direction * y as f32 * pixel_height;
+                let ray = Ray {
+                    origin: tree_view.spyglass.viewport.origin,
+                    direction: (glass_point - tree_view.spyglass.viewport.origin).normalized(),
+                };
+
+                use std::io::Write;
+                std::io::stdout().flush().ok().unwrap();
+
+                if let Some(hit) = tree.tree.get_by_ray(&ray) {
+                    let (data, _, normal) = hit;
+                    //Because both vector should be normalized, the dot product should be 1*1*cos(angle)
+                    //That means it is in range -1, +1, which should be accounted for
+                    let diffuse_light_strength =
+                        1. - (normal.dot(&diffuse_light_normal) / 2. + 0.5);
+                    img.put_pixel(
+                        x,
+                        actual_y_in_image,
+                        Rgb([
+                            (data.r as f32 * diffuse_light_strength) as u8,
+                            (data.g as f32 * diffuse_light_strength) as u8,
+                            (data.b as f32 * diffuse_light_strength) as u8,
+                        ]),
+                    );
+                } else {
+                    img.put_pixel(x, actual_y_in_image, Rgb([128, 128, 128]));
+                }
+            }
+        }
+
+        img.save("example_junk_cpu_render.png").ok().unwrap();
+    }
+
     let multiplier = if keys.pressed(KeyCode::ShiftLeft) {
         10.0 // Doesn't have any effect?!
     } else {
         1.0
     };
+
     if keys.pressed(KeyCode::ArrowUp) {
         angles_query.single_mut().roll = angle_update_fn(angles_query.single().roll, ADDITION);
     }
@@ -209,11 +254,9 @@ fn handle_zoom(
     }
     if keys.pressed(KeyCode::ArrowLeft) {
         angles_query.single_mut().yaw = angle_update_fn(angles_query.single().yaw, ADDITION);
-        // println!("viewport: {:?}", viewing_glass.viewport);
     }
     if keys.pressed(KeyCode::ArrowRight) {
         angles_query.single_mut().yaw = angle_update_fn(angles_query.single().yaw, -ADDITION);
-        // println!("viewport: {:?}", viewing_glass.viewport);
     }
     if keys.pressed(KeyCode::PageUp) {
         angles_query.single_mut().radius *= 1. - 0.02 * multiplier;
@@ -222,12 +265,12 @@ fn handle_zoom(
         angles_query.single_mut().radius *= 1. + 0.02 * multiplier;
     }
     if keys.pressed(KeyCode::Home) {
-        viewing_glass.viewport.w_h_fov.x *= 1. + 0.09 * multiplier;
-        viewing_glass.viewport.w_h_fov.y *= 1. + 0.09 * multiplier;
+        tree_view.spyglass.viewport.w_h_fov.x *= 1. + 0.09 * multiplier;
+        tree_view.spyglass.viewport.w_h_fov.y *= 1. + 0.09 * multiplier;
     }
     if keys.pressed(KeyCode::End) {
-        viewing_glass.viewport.w_h_fov.x *= 1. - 0.09 * multiplier;
-        viewing_glass.viewport.w_h_fov.y *= 1. - 0.09 * multiplier;
+        tree_view.spyglass.viewport.w_h_fov.x *= 1. - 0.09 * multiplier;
+        tree_view.spyglass.viewport.w_h_fov.y *= 1. - 0.09 * multiplier;
     }
 }
 
diff --git a/examples/minecraft.rs b/examples/minecraft.rs
index d4ce042..da49963 100644
--- a/examples/minecraft.rs
+++ b/examples/minecraft.rs
@@ -1,15 +1,10 @@
-#[cfg(feature = "bevy_wgpu")]
-use shocovox_rs::octree::Octree;
-
 #[cfg(feature = "bevy_wgpu")]
 use bevy::{prelude::*, window::WindowPlugin};
 
 #[cfg(feature = "bevy_wgpu")]
 use shocovox_rs::octree::{
-    raytracing::{
-        bevy::create_viewing_glass, Ray, ShocoVoxRenderPlugin, ShocoVoxViewingGlass, Viewport,
-    },
-    Albedo, V3c, VoxelData,
+    raytracing::{OctreeGPUHost, Ray, SvxViewSet, Viewport},
+    Albedo, Octree, V3c,
 };
 
 #[cfg(feature = "bevy_wgpu")]
@@ -25,15 +20,6 @@ const DISPLAY_RESOLUTION: [u32; 2] = [1024, 768];
 #[cfg(feature = "bevy_wgpu")]
 const BRICK_DIMENSION: usize = 32;
 
-#[cfg(feature = "bevy_wgpu")]
-#[derive(Resource)]
-struct TreeResource<T, const DIM: usize>
-where
-    T: Default + Clone + PartialEq + VoxelData,
-{
-    tree: Octree<T, DIM>,
-}
-
 #[cfg(feature = "bevy_wgpu")]
 fn main() {
     App::new()
@@ -47,9 +33,9 @@ fn main() {
                 }),
                 ..default()
             }),
-            ShocoVoxRenderPlugin {
-                resolution: DISPLAY_RESOLUTION,
-            },
+            shocovox_rs::octree::raytracing::RenderBevyPlugin::<Albedo, BRICK_DIMENSION>::new(
+                DISPLAY_RESOLUTION,
+            ),
             bevy::diagnostic::FrameTimeDiagnosticsPlugin,
             PerfUiPlugin,
         ))
@@ -78,21 +64,23 @@ 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.,
-        radius: tree.get_size() as f32 * 2.2,
+        radius: tree.get_size() as f32 * 0.8,
     });
 
-    let render_data = tree.create_bevy_view();
-    let viewing_glass = create_viewing_glass(
-        &Viewport {
-            origin,
+    let mut host = OctreeGPUHost { tree };
+    let mut views = SvxViewSet::default();
+    let output_texture = host.create_new_view(
+        &mut views,
+        35,
+        Viewport {
+            origin: V3c {
+                x: 0.,
+                y: 0.,
+                z: 0.,
+            },
             direction: V3c {
                 x: 0.,
                 y: 0.,
@@ -103,19 +91,17 @@ fn setup(mut commands: Commands, images: ResMut<Assets<Image>>) {
         DISPLAY_RESOLUTION,
         images,
     );
+    commands.insert_resource(host);
+    commands.insert_resource(views);
     commands.spawn(SpriteBundle {
         sprite: Sprite {
             custom_size: Some(Vec2::new(1024., 768.)),
             ..default()
         },
-        texture: viewing_glass.output_texture.clone(),
+        texture: output_texture,
         ..default()
     });
     commands.spawn(Camera2dBundle::default());
-    commands.insert_resource(TreeResource { tree });
-    commands.insert_resource(render_data);
-    commands.insert_resource(viewing_glass);
-
     commands.spawn((
         PerfUiRoot::default(),
         PerfUiEntryFPS {
@@ -144,29 +130,27 @@ struct DomePosition {
 }
 
 #[cfg(feature = "bevy_wgpu")]
-fn rotate_camera(
-    angles_query: Query<&mut DomePosition>,
-    mut viewing_glass: ResMut<ShocoVoxViewingGlass>,
-) {
+fn rotate_camera(angles_query: Query<&mut DomePosition>, view_set: ResMut<SvxViewSet>) {
     let (yaw, roll) = (angles_query.single().yaw, angles_query.single().roll);
-
     let radius = angles_query.single().radius;
-    viewing_glass.viewport.origin = V3c::new(
+    let mut tree_view = view_set.views[0].lock().unwrap();
+    tree_view.spyglass.viewport.origin = V3c::new(
         radius / 2. + yaw.sin() * radius,
         radius + roll.sin() * radius * 2.,
         radius / 2. + yaw.cos() * radius,
     );
-    viewing_glass.viewport.direction =
-        (V3c::unit(radius / 2.) - viewing_glass.viewport.origin).normalized();
+    tree_view.spyglass.viewport.direction =
+        (V3c::unit(radius / 2.) - tree_view.spyglass.viewport.origin).normalized();
 }
 
 #[cfg(feature = "bevy_wgpu")]
 fn handle_zoom(
     keys: Res<ButtonInput<KeyCode>>,
-    mut viewing_glass: ResMut<ShocoVoxViewingGlass>,
+    tree: ResMut<OctreeGPUHost<Albedo, BRICK_DIMENSION>>,
+    view_set: ResMut<SvxViewSet>,
     mut angles_query: Query<&mut DomePosition>,
-    tree: Res<TreeResource<Albedo, BRICK_DIMENSION>>,
 ) {
+    let mut tree_view = view_set.views[0].lock().unwrap();
     const ADDITION: f32 = 0.05;
     let angle_update_fn = |angle, delta| -> f32 {
         let new_angle = angle + delta;
@@ -180,14 +164,16 @@ fn handle_zoom(
         // Render the current view with CPU
         let viewport_up_direction = V3c::new(0., 1., 0.);
         let viewport_right_direction = viewport_up_direction
-            .cross(viewing_glass.viewport.direction)
+            .cross(tree_view.spyglass.viewport.direction)
             .normalized();
-        let pixel_width = viewing_glass.viewport.w_h_fov.x as f32 / DISPLAY_RESOLUTION[0] as f32;
-        let pixel_height = viewing_glass.viewport.w_h_fov.y as f32 / DISPLAY_RESOLUTION[1] as f32;
-        let viewport_bottom_left = viewing_glass.viewport.origin
-            + (viewing_glass.viewport.direction * viewing_glass.viewport.w_h_fov.z)
-            - (viewport_up_direction * (viewing_glass.viewport.w_h_fov.y / 2.))
-            - (viewport_right_direction * (viewing_glass.viewport.w_h_fov.x / 2.));
+        let pixel_width =
+            tree_view.spyglass.viewport.w_h_fov.x as f32 / DISPLAY_RESOLUTION[0] as f32;
+        let pixel_height =
+            tree_view.spyglass.viewport.w_h_fov.y as f32 / DISPLAY_RESOLUTION[1] as f32;
+        let viewport_bottom_left = tree_view.spyglass.viewport.origin
+            + (tree_view.spyglass.viewport.direction * tree_view.spyglass.viewport.w_h_fov.z)
+            - (viewport_up_direction * (tree_view.spyglass.viewport.w_h_fov.y / 2.))
+            - (viewport_right_direction * (tree_view.spyglass.viewport.w_h_fov.x / 2.));
 
         // define light
         let diffuse_light_normal = V3c::new(0., -1., 1.).normalized();
@@ -205,8 +191,8 @@ fn handle_zoom(
                     + viewport_right_direction * x as f32 * pixel_width
                     + viewport_up_direction * y as f32 * pixel_height;
                 let ray = Ray {
-                    origin: viewing_glass.viewport.origin,
-                    direction: (glass_point - viewing_glass.viewport.origin).normalized(),
+                    origin: tree_view.spyglass.viewport.origin,
+                    direction: (glass_point - tree_view.spyglass.viewport.origin).normalized(),
                 };
 
                 use std::io::Write;
@@ -250,11 +236,9 @@ fn handle_zoom(
     }
     if keys.pressed(KeyCode::ArrowLeft) {
         angles_query.single_mut().yaw = angle_update_fn(angles_query.single().yaw, ADDITION);
-        // println!("viewport: {:?}", viewing_glass.viewport);
     }
     if keys.pressed(KeyCode::ArrowRight) {
         angles_query.single_mut().yaw = angle_update_fn(angles_query.single().yaw, -ADDITION);
-        // println!("viewport: {:?}", viewing_glass.viewport);
     }
     if keys.pressed(KeyCode::PageUp) {
         angles_query.single_mut().radius *= 1. - 0.02 * multiplier;
@@ -263,12 +247,12 @@ fn handle_zoom(
         angles_query.single_mut().radius *= 1. + 0.02 * multiplier;
     }
     if keys.pressed(KeyCode::Home) {
-        viewing_glass.viewport.w_h_fov.x *= 1. + 0.09 * multiplier;
-        viewing_glass.viewport.w_h_fov.y *= 1. + 0.09 * multiplier;
+        tree_view.spyglass.viewport.w_h_fov.x *= 1. + 0.09 * multiplier;
+        tree_view.spyglass.viewport.w_h_fov.y *= 1. + 0.09 * multiplier;
     }
     if keys.pressed(KeyCode::End) {
-        viewing_glass.viewport.w_h_fov.x *= 1. - 0.09 * multiplier;
-        viewing_glass.viewport.w_h_fov.y *= 1. - 0.09 * multiplier;
+        tree_view.spyglass.viewport.w_h_fov.x *= 1. - 0.09 * multiplier;
+        tree_view.spyglass.viewport.w_h_fov.y *= 1. - 0.09 * multiplier;
     }
 }
 
diff --git a/src/octree/mod.rs b/src/octree/mod.rs
index 2f3c791..0efe2d3 100644
--- a/src/octree/mod.rs
+++ b/src/octree/mod.rs
@@ -92,7 +92,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;
@@ -248,7 +248,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;
diff --git a/src/octree/node.rs b/src/octree/node.rs
index 059aed2..0153528 100644
--- a/src/octree/node.rs
+++ b/src/octree/node.rs
@@ -7,9 +7,24 @@ use crate::spatial::{
     math::{set_occupancy_in_bitmap_64bits, BITMAP_DIMENSION},
 };
 
-///####################################################################################
-/// NodeChildren
-///####################################################################################
+//####################################################################################
+//  ██████   █████    ███████    ██████████   ██████████
+// ░░██████ ░░███   ███░░░░░███ ░░███░░░░███ ░░███░░░░░█
+//  ░███░███ ░███  ███     ░░███ ░███   ░░███ ░███  █ ░
+//  ░███░░███░███ ░███      ░███ ░███    ░███ ░██████
+//  ░███ ░░██████ ░███      ░███ ░███    ░███ ░███░░█
+//  ░███  ░░█████ ░░███     ███  ░███    ███  ░███ ░   █
+//  █████  ░░█████ ░░░███████░   ██████████   ██████████
+// ░░░░░    ░░░░░    ░░░░░░░    ░░░░░░░░░░   ░░░░░░░░░░
+//    █████████  █████   █████ █████ █████       ██████████   ███████████   ██████████ ██████   █████
+//   ███░░░░░███░░███   ░░███ ░░███ ░░███       ░░███░░░░███ ░░███░░░░░███ ░░███░░░░░█░░██████ ░░███
+//  ███     ░░░  ░███    ░███  ░███  ░███        ░███   ░░███ ░███    ░███  ░███  █ ░  ░███░███ ░███
+// ░███          ░███████████  ░███  ░███        ░███    ░███ ░██████████   ░██████    ░███░░███░███
+// ░███          ░███░░░░░███  ░███  ░███        ░███    ░███ ░███░░░░░███  ░███░░█    ░███ ░░██████
+// ░░███     ███ ░███    ░███  ░███  ░███      █ ░███    ███  ░███    ░███  ░███ ░   █ ░███  ░░█████
+//  ░░█████████  █████   █████ █████ ███████████ ██████████   █████   █████ ██████████ █████  ░░█████
+//   ░░░░░░░░░  ░░░░░   ░░░░░ ░░░░░ ░░░░░░░░░░░ ░░░░░░░░░░   ░░░░░   ░░░░░ ░░░░░░░░░░ ░░░░░    ░░░░░
+//####################################################################################
 impl<T> NodeChildren<T>
 where
     T: Default + Clone + Eq,
@@ -85,12 +100,27 @@ where
     }
 }
 
-///####################################################################################
-/// BrickData
-///####################################################################################
+//####################################################################################
+//  ███████████  ███████████   █████   █████████  █████   ████
+// ░░███░░░░░███░░███░░░░░███ ░░███   ███░░░░░███░░███   ███░
+//  ░███    ░███ ░███    ░███  ░███  ███     ░░░  ░███  ███
+//  ░██████████  ░██████████   ░███ ░███          ░███████
+//  ░███░░░░░███ ░███░░░░░███  ░███ ░███          ░███░░███
+//  ░███    ░███ ░███    ░███  ░███ ░░███     ███ ░███ ░░███
+//  ███████████  █████   █████ █████ ░░█████████  █████ ░░████
+// ░░░░░░░░░░░  ░░░░░   ░░░░░ ░░░░░   ░░░░░░░░░  ░░░░░   ░░░░
+//  ██████████     █████████   ███████████   █████████
+// ░░███░░░░███   ███░░░░░███ ░█░░░███░░░█  ███░░░░░███
+//  ░███   ░░███ ░███    ░███ ░   ░███  ░  ░███    ░███
+//  ░███    ░███ ░███████████     ░███     ░███████████
+//  ░███    ░███ ░███░░░░░███     ░███     ░███░░░░░███
+//  ░███    ███  ░███    ░███     ░███     ░███    ░███
+//  ██████████   █████   █████    █████    █████   █████
+// ░░░░░░░░░░   ░░░░░   ░░░░░    ░░░░░    ░░░░░   ░░░░░
+//####################################################################################
 impl<T, const DIM: usize> BrickData<T, DIM>
 where
-    T: VoxelData + PartialEq + Clone + Copy + Default,
+    T: VoxelData + PartialEq + Clone + Default,
 {
     /// Provides occupancy information for the part of the brick corresponmding
     /// to the given octant based on the contents of the brick
@@ -231,7 +261,7 @@ where
             if homogeneous_type.is_empty() {
                 *self = BrickData::Empty;
             } else {
-                *self = BrickData::Solid(*homogeneous_type);
+                *self = BrickData::Solid(homogeneous_type.clone());
             }
             true
         } else {
@@ -240,9 +270,24 @@ where
     }
 }
 
-///####################################################################################
-/// NodeContent
-///####################################################################################
+//####################################################################################
+//  ██████   █████    ███████    ██████████   ██████████
+// ░░██████ ░░███   ███░░░░░███ ░░███░░░░███ ░░███░░░░░█
+//  ░███░███ ░███  ███     ░░███ ░███   ░░███ ░███  █ ░
+//  ░███░░███░███ ░███      ░███ ░███    ░███ ░██████
+//  ░███ ░░██████ ░███      ░███ ░███    ░███ ░███░░█
+//  ░███  ░░█████ ░░███     ███  ░███    ███  ░███ ░   █
+//  █████  ░░█████ ░░░███████░   ██████████   ██████████
+// ░░░░░    ░░░░░    ░░░░░░░    ░░░░░░░░░░   ░░░░░░░░░░
+//    █████████     ███████    ██████   █████ ███████████ ██████████ ██████   █████ ███████████
+//   ███░░░░░███  ███░░░░░███ ░░██████ ░░███ ░█░░░███░░░█░░███░░░░░█░░██████ ░░███ ░█░░░███░░░█
+//  ███     ░░░  ███     ░░███ ░███░███ ░███ ░   ░███  ░  ░███  █ ░  ░███░███ ░███ ░   ░███  ░
+// ░███         ░███      ░███ ░███░░███░███     ░███     ░██████    ░███░░███░███     ░███
+// ░███         ░███      ░███ ░███ ░░██████     ░███     ░███░░█    ░███ ░░██████     ░███
+// ░░███     ███░░███     ███  ░███  ░░█████     ░███     ░███ ░   █ ░███  ░░█████     ░███
+//  ░░█████████  ░░░███████░   █████  ░░█████    █████    ██████████ █████  ░░█████    █████
+//   ░░░░░░░░░     ░░░░░░░    ░░░░░    ░░░░░    ░░░░░    ░░░░░░░░░░ ░░░░░    ░░░░░    ░░░░░
+//####################################################################################
 
 impl<T, const DIM: usize> NodeContent<T, DIM>
 where
diff --git a/src/octree/raytracing/bevy/cache.rs b/src/octree/raytracing/bevy/cache.rs
new file mode 100644
index 0000000..327126c
--- /dev/null
+++ b/src/octree/raytracing/bevy/cache.rs
@@ -0,0 +1,690 @@
+use crate::object_pool::empty_marker;
+use crate::octree::raytracing::bevy::types::BrickOwnedBy;
+use crate::spatial::math::flat_projection;
+use crate::{
+    octree::{
+        raytracing::bevy::types::{OctreeRenderData, Voxelement},
+        types::{NodeChildrenArray, NodeContent},
+        BrickData, Octree, VoxelData,
+    },
+    spatial::lut::BITMAP_MASK_FOR_OCTANT_LUT,
+};
+use bevy::math::Vec4;
+
+use super::types::{OctreeGPUDataHandler, VictimPointer};
+
+//##############################################################################
+//  █████   █████ █████   █████████  ███████████ █████ ██████   ██████
+// ░░███   ░░███ ░░███   ███░░░░░███░█░░░███░░░█░░███ ░░██████ ██████
+//  ░███    ░███  ░███  ███     ░░░ ░   ░███  ░  ░███  ░███░█████░███
+//  ░███    ░███  ░███ ░███             ░███     ░███  ░███░░███ ░███
+//  ░░███   ███   ░███ ░███             ░███     ░███  ░███ ░░░  ░███
+//   ░░░█████░    ░███ ░░███     ███    ░███     ░███  ░███      ░███
+//     ░░███      █████ ░░█████████     █████    █████ █████     █████
+//      ░░░      ░░░░░   ░░░░░░░░░     ░░░░░    ░░░░░ ░░░░░     ░░░░░
+//  ███████████  ███████████ ███████████
+// ░░███░░░░░███░█░░░███░░░█░░███░░░░░███
+//  ░███    ░███░   ░███  ░  ░███    ░███
+//  ░██████████     ░███     ░██████████
+//  ░███░░░░░░      ░███     ░███░░░░░███
+//  ░███            ░███     ░███    ░███
+//  █████           █████    █████   █████
+// ░░░░░           ░░░░░    ░░░░░   ░░░░░
+//##############################################################################
+impl VictimPointer {
+    /// Returns the size of the buffer this pointer covers
+    pub(crate) fn len(&self) -> usize {
+        self.max_meta_len
+    }
+
+    /// Returns true if no new nodes can be added without overwriting another
+    pub(crate) fn is_full(&self) -> bool {
+        self.max_meta_len <= self.stored_items
+    }
+
+    /// Creates object, based on the given cache length it should cover
+    pub(crate) fn new(max_meta_len: usize) -> Self {
+        Self {
+            max_meta_len,
+            loop_count: 0,
+            stored_items: 0,
+            meta_index: max_meta_len - 1,
+            child: 0,
+        }
+    }
+
+    /// Steps the iterator forward to the next children, if available, or the next node.
+    /// Wraps around
+    pub(crate) fn step(&mut self) {
+        if self.child >= 7 {
+            self.skip_node();
+        } else {
+            self.child += 1;
+        }
+    }
+
+    /// Steps the iterator forward one node
+    pub(crate) fn skip_node(&mut self) {
+        if self.meta_index == 0 {
+            self.loop_count += 1;
+            self.meta_index = self.max_meta_len - 1;
+        } else {
+            self.meta_index -= 1;
+        }
+        self.child = 0;
+    }
+
+    /// Provides the number of times the victim node pointer has started from the first element in the cache
+    pub(crate) fn get_loop_count(&self) -> usize {
+        self.loop_count
+    }
+
+    /// Provides the first available index in the metadata buffer which can be overwritten
+    /// with node related meta information and optionally the source where the child was taken from.
+    fn first_available_node(
+        &mut self,
+        render_data: &mut OctreeRenderData,
+    ) -> (usize, Option<(usize, u8)>) {
+        // If there is space left in the cache, use it all up
+        if !self.is_full() {
+            render_data.metadata[self.stored_items] |= OctreeGPUDataHandler::NODE_USED_MASK;
+            self.meta_index = self.stored_items;
+            self.stored_items += 1;
+            return (self.meta_index, None);
+        }
+
+        //look for the next internal node ( with node children )
+        loop {
+            // child at target is not empty in a non-leaf node, which means
+            // the target child might point to an internal node if it's valid
+            // parent node has a child at target octant, which isn't invalid
+            if 0 == (render_data.metadata[self.meta_index] & OctreeGPUDataHandler::NODE_LEAF_MASK)
+                && render_data.node_children[self.meta_index * 8 + self.child] != empty_marker()
+            {
+                let child_meta_index =
+                    render_data.node_children[self.meta_index * 8 + self.child] as usize;
+                if 0 == (render_data.metadata[child_meta_index]
+                    & OctreeGPUDataHandler::NODE_USED_MASK)
+                {
+                    render_data.metadata[child_meta_index] |= OctreeGPUDataHandler::NODE_USED_MASK;
+                    return (child_meta_index, Some((self.meta_index, self.child as u8)));
+                } else {
+                    // mark child as unused
+                    render_data.metadata[child_meta_index] &= !OctreeGPUDataHandler::NODE_USED_MASK;
+                }
+            }
+            self.step();
+        }
+    }
+}
+
+impl OctreeGPUDataHandler {
+    //##############################################################################
+    //  ██████████     █████████   ███████████   █████████
+    // ░░███░░░░███   ███░░░░░███ ░█░░░███░░░█  ███░░░░░███
+    //  ░███   ░░███ ░███    ░███ ░   ░███  ░  ░███    ░███
+    //  ░███    ░███ ░███████████     ░███     ░███████████
+    //  ░███    ░███ ░███░░░░░███     ░███     ░███░░░░░███
+    //  ░███    ███  ░███    ░███     ░███     ░███    ░███
+    //  ██████████   █████   █████    █████    █████   █████
+    // ░░░░░░░░░░   ░░░░░   ░░░░░    ░░░░░    ░░░░░   ░░░░░
+
+    //  ██████████   ██████████  █████████  █████   █████████  ██████   █████
+    // ░░███░░░░███ ░░███░░░░░█ ███░░░░░███░░███   ███░░░░░███░░██████ ░░███
+    //  ░███   ░░███ ░███  █ ░ ░███    ░░░  ░███  ███     ░░░  ░███░███ ░███
+    //  ░███    ░███ ░██████   ░░█████████  ░███ ░███          ░███░░███░███
+    //  ░███    ░███ ░███░░█    ░░░░░░░░███ ░███ ░███    █████ ░███ ░░██████
+    //  ░███    ███  ░███ ░   █ ███    ░███ ░███ ░░███  ░░███  ░███  ░░█████
+    //  ██████████   ██████████░░█████████  █████ ░░█████████  █████  ░░█████
+    // ░░░░░░░░░░   ░░░░░░░░░░  ░░░░░░░░░  ░░░░░   ░░░░░░░░░  ░░░░░    ░░░░░
+    //##############################################################################
+
+    /// Bitmask in metadata where the non-zero bits represent if the given node is used
+    const NODE_USED_MASK: u32 = 0x00000001;
+
+    /// Bitmask in metadata where the non-zero bits represent if the given node is a leaf
+    const NODE_LEAF_MASK: u32 = 0x00000004;
+
+    /// Bitmask in metadata where the non-zero bits represent if the given leaf is uniform
+    /// Note: Non-leaf nodes can't be uniform
+    const NODE_UNIFORM_MASK: u32 = 0x00000008;
+
+    /// Provides the mask used with one metadata element to signal that the contained brick is used.
+    /// Index of the metadata element should be brick index divided by 8, as one metadata element contains 8 bricks
+    fn brick_used_mask(brick_index: usize) -> u32 {
+        0x01 << (24 + (brick_index % 8))
+    }
+
+    /// Updates the meta element value to store the brick structure of the given leaf node.
+    /// Does not erase anything in @sized_node_meta, it's expected to be cleared before
+    /// the first use of this function
+    /// for the given brick octant
+    /// * `sized_node_meta` - the bytes to update
+    /// * `brick` - the brick to describe into the bytes
+    /// * `brick_octant` - the octant to update in the bytes
+    fn meta_add_leaf_brick_structure<T, const DIM: usize>(
+        sized_node_meta: &mut u32,
+        brick: &BrickData<T, DIM>,
+        brick_octant: usize,
+    ) where
+        T: Default + Clone + PartialEq + VoxelData,
+    {
+        match brick {
+            BrickData::Empty => {} // Child structure properties already set to NIL
+            BrickData::Solid(_voxel) => {
+                // set child Occupied bits, child Structure bits already set to NIL
+                *sized_node_meta |= 0x01 << (8 + brick_octant);
+            }
+            BrickData::Parted(_brick) => {
+                // set child Occupied bits
+                *sized_node_meta |= 0x01 << (8 + brick_octant);
+
+                // set child Structure bits
+                *sized_node_meta |= 0x01 << (16 + brick_octant);
+            }
+        };
+    }
+
+    /// Creates the descriptor bytes for the given node
+    fn create_node_properties<T, const DIM: usize>(node: &NodeContent<T, DIM>) -> u32
+    where
+        T: Default + Copy + Clone + PartialEq + VoxelData,
+    {
+        let mut meta = 0;
+        match node {
+            NodeContent::Internal(_) | NodeContent::Nothing => {
+                meta &= !Self::NODE_LEAF_MASK; // element is not leaf
+                meta &= !Self::NODE_UNIFORM_MASK; // element is not uniform
+            }
+            NodeContent::Leaf(bricks) => {
+                meta |= Self::NODE_LEAF_MASK; // element is leaf
+                meta &= !Self::NODE_UNIFORM_MASK; // element is not uniform
+                for octant in 0..8 {
+                    Self::meta_add_leaf_brick_structure(&mut meta, &bricks[octant], octant);
+                }
+            }
+            NodeContent::UniformLeaf(brick) => {
+                meta |= Self::NODE_LEAF_MASK; // element is leaf
+                meta |= Self::NODE_UNIFORM_MASK; // element is uniform
+                Self::meta_add_leaf_brick_structure(&mut meta, brick, 0);
+            }
+        };
+        meta
+    }
+
+    //##############################################################################
+    //  ██████████ ███████████     █████████    █████████  ██████████
+    // ░░███░░░░░█░░███░░░░░███   ███░░░░░███  ███░░░░░███░░███░░░░░█
+    //  ░███  █ ░  ░███    ░███  ░███    ░███ ░███    ░░░  ░███  █ ░
+    //  ░██████    ░█████████    ░███████████ ░░█████████  ░██████
+    //  ░███░░█    ░███░░░░░███  ░███░░░░░███  ░░░░░░░░███ ░███░░█
+    //  ░███ ░   █ ░███    ░███  ░███    ░███  ███    ░███ ░███ ░   █
+    //  ██████████ █████   █████ █████   █████░░█████████  ██████████
+    // ░░░░░░░░░░ ░░░░░   ░░░░░ ░░░░░   ░░░░░  ░░░░░░░░░  ░░░░░░░░░░
+    //  ██████   █████    ███████    ██████████   ██████████
+    // ░░██████ ░░███   ███░░░░░███ ░░███░░░░███ ░░███░░░░░█
+    //  ░███░███ ░███  ███     ░░███ ░███   ░░███ ░███  █ ░
+    //  ░███░░███░███ ░███      ░███ ░███    ░███ ░██████
+    //  ░███ ░░██████ ░███      ░███ ░███    ░███ ░███░░█
+    //  ░███  ░░█████ ░░███     ███  ░███    ███  ░███ ░   █
+    //  █████  ░░█████ ░░░███████░   ██████████   ██████████
+    // ░░░░░    ░░░░░    ░░░░░░░    ░░░░░░░░░░   ░░░░░░░░░░
+    //    █████████  █████   █████ █████ █████       ██████████
+    //   ███░░░░░███░░███   ░░███ ░░███ ░░███       ░░███░░░░███
+    //  ███     ░░░  ░███    ░███  ░███  ░███        ░███   ░░███
+    // ░███          ░███████████  ░███  ░███        ░███    ░███
+    // ░███          ░███░░░░░███  ░███  ░███        ░███    ░███
+    // ░░███     ███ ░███    ░███  ░███  ░███      █ ░███    ███
+    //  ░░█████████  █████   █████ █████ ███████████ ██████████
+    //   ░░░░░░░░░  ░░░░░   ░░░░░ ░░░░░ ░░░░░░░░░░░ ░░░░░░░░░░
+    //##############################################################################
+    /// Erases the child node pointed by the given victim pointer
+    /// returns with the vector of node index values and brick index values modified
+    fn erase_node_child<T, const DIM: usize>(
+        &mut self,
+        meta_index: usize,
+        child_octant: usize,
+        tree: &Octree<T, DIM>,
+    ) -> (Vec<usize>, Vec<usize>)
+    where
+        T: Default + Clone + PartialEq + VoxelData,
+    {
+        let mut modified_bricks = Vec::new();
+        let mut modified_nodes = vec![meta_index];
+        debug_assert!(
+            self.node_key_vs_meta_index.contains_right(&meta_index),
+            "Expected parent node to be in metadata index hash! (meta: {meta_index})"
+        );
+        let parent_key = self
+            .node_key_vs_meta_index
+            .get_by_right(&meta_index)
+            .unwrap();
+
+        debug_assert!(
+            tree.nodes.key_is_valid(*parent_key),
+            "Expected parent node({:?}) to be valid",
+            parent_key
+        );
+
+        // Erase connection to parent
+        let child_index = self.render_data.node_children[meta_index * 8 + child_octant] as usize;
+        self.render_data.node_children[meta_index * 8 + child_octant] = empty_marker();
+        debug_assert_ne!(
+            child_index,
+            empty_marker() as usize,
+            "Expected victim pointer to point to an erasable node/brick, instead of: {child_index}"
+        );
+
+        match tree.nodes.get(*parent_key) {
+            NodeContent::Nothing => {
+                panic!("HOW DO I ERASE NOTHING. AMERICA EXPLAIN")
+            }
+            NodeContent::Internal(_occupied_bits) => {
+                debug_assert!(
+                    self.node_key_vs_meta_index.contains_right(&child_index),
+                    "Expected erased child node index[{child_index}] to be in metadata index hash!"
+                );
+                let child_key = self
+                    .node_key_vs_meta_index
+                    .get_by_right(&child_index)
+                    .unwrap();
+                debug_assert!(
+                    tree.nodes.key_is_valid(*child_key),
+                    "Expected erased child node({child_key}) to be valid"
+                );
+
+                if let NodeContent::Leaf(_) | NodeContent::UniformLeaf(_) =
+                    tree.nodes.get(*child_key)
+                {
+                    // make the children bricks of the removed leaf orphan
+                    for octant in 0..8 {
+                        let brick_index =
+                            self.render_data.node_children[child_index * 8 + octant] as usize;
+                        if brick_index != empty_marker() as usize {
+                            self.brick_ownership[brick_index] = BrickOwnedBy::NotOwned;
+
+                            // mark brick as unused
+                            self.render_data.metadata[brick_index / 8] &=
+                                !Self::brick_used_mask(brick_index);
+
+                            // Eliminate connection
+                            self.render_data.node_children[child_index * 8 + octant] =
+                                empty_marker();
+
+                            modified_bricks.push(brick_index);
+                        }
+                    }
+                }
+                modified_nodes.push(child_index);
+            }
+            NodeContent::UniformLeaf(_) | NodeContent::Leaf(_) => {
+                debug_assert!(
+                    (0 == child_octant)
+                        || matches!(tree.nodes.get(*parent_key), NodeContent::Leaf(_)),
+                    "Expected child octant in uniform leaf to be 0 in: {:?}",
+                    (meta_index, child_octant)
+                );
+                if child_index != empty_marker() as usize {
+                    self.brick_ownership[child_index] = BrickOwnedBy::NotOwned;
+                    modified_bricks.push(child_index);
+
+                    // mark brick as unused
+                    self.render_data.metadata[child_index / 8] &=
+                        !Self::brick_used_mask(child_index);
+                }
+            }
+        }
+
+        //return with updated ranges in voxels and metadata
+        (modified_nodes, modified_bricks)
+    }
+
+    //##############################################################################
+    //    █████████   ██████████   ██████████
+    //   ███░░░░░███ ░░███░░░░███ ░░███░░░░███
+    //  ░███    ░███  ░███   ░░███ ░███   ░░███
+    //  ░███████████  ░███    ░███ ░███    ░███
+    //  ░███░░░░░███  ░███    ░███ ░███    ░███
+    //  ░███    ░███  ░███    ███  ░███    ███
+    //  █████   █████ ██████████   ██████████
+    // ░░░░░   ░░░░░ ░░░░░░░░░░   ░░░░░░░░░░
+    //  ██████   █████    ███████    ██████████   ██████████
+    // ░░██████ ░░███   ███░░░░░███ ░░███░░░░███ ░░███░░░░░█
+    //  ░███░███ ░███  ███     ░░███ ░███   ░░███ ░███  █ ░
+    //  ░███░░███░███ ░███      ░███ ░███    ░███ ░██████
+    //  ░███ ░░██████ ░███      ░███ ░███    ░███ ░███░░█
+    //  ░███  ░░█████ ░░███     ███  ░███    ███  ░███ ░   █
+    //  █████  ░░█████ ░░░███████░   ██████████   ██████████
+    // ░░░░░    ░░░░░    ░░░░░░░    ░░░░░░░░░░   ░░░░░░░░░░
+    //##############################################################################
+    /// Writes the data of the node to the first available index
+    /// Upon success, returns the index in metadata where the node is added
+    /// and vectors of modified nodes, bricks:
+    /// (meta_index, modified_nodes, modified_bricks)
+    pub(crate) fn add_node<T, const DIM: usize>(
+        &mut self,
+        tree: &Octree<T, DIM>,
+        node_key: usize,
+        try_add_children: bool,
+    ) -> Option<(usize, Vec<usize>, Vec<usize>)>
+    where
+        T: Default + Copy + Clone + PartialEq + VoxelData + Send + Sync + 'static,
+    {
+        if try_add_children && self.victim_node.is_full() {
+            // Do not add additional nodes at initial upload if the cache is already full
+            return None;
+        }
+
+        // Determine the index in meta, overwrite a currently present node if needed
+        let (node_element_index, robbed_parent) =
+            self.victim_node.first_available_node(&mut self.render_data);
+        let (mut modified_nodes, mut modified_bricks) = if let Some(robbed_parent) = robbed_parent {
+            self.erase_node_child(robbed_parent.0, robbed_parent.1 as usize, tree)
+        } else {
+            (vec![node_element_index], Vec::new())
+        };
+
+        self.node_key_vs_meta_index
+            .insert(node_key, node_element_index);
+
+        // Add node properties to metadata
+        self.render_data.metadata[node_element_index] =
+            Self::create_node_properties(tree.nodes.get(node_key));
+
+        // Update occupancy in ocbits
+        let occupied_bits = tree.stored_occupied_bits(node_key);
+        self.render_data.node_ocbits[node_element_index * 2] =
+            (occupied_bits & 0x00000000FFFFFFFF) as u32;
+        self.render_data.node_ocbits[node_element_index * 2 + 1] =
+            ((occupied_bits & 0xFFFFFFFF00000000) >> 32) as u32;
+
+        // Add node content
+        match tree.nodes.get(node_key) {
+            NodeContent::UniformLeaf(brick) => {
+                debug_assert!(
+                    matches!(
+                        tree.node_children[node_key].content,
+                        NodeChildrenArray::OccupancyBitmap(_)
+                    ),
+                    "Expected Uniform leaf to have OccupancyBitmap(_) instead of {:?}",
+                    tree.node_children[node_key].content
+                );
+
+                if try_add_children {
+                    let (brick_index, mut current_modified_nodes, mut current_modified_bricks) =
+                        self.add_brick(tree, node_key, 0);
+                    modified_bricks.push(brick_index as usize);
+                    modified_nodes.append(&mut current_modified_nodes);
+                    modified_bricks.append(&mut current_modified_bricks);
+
+                    self.render_data.node_children[node_element_index * 8 + 0] = brick_index;
+                } else {
+                    self.render_data.node_children[node_element_index * 8 + 0] = empty_marker();
+                }
+
+                self.render_data.node_children[node_element_index * 8 + 1] = empty_marker();
+                self.render_data.node_children[node_element_index * 8 + 2] = empty_marker();
+                self.render_data.node_children[node_element_index * 8 + 3] = empty_marker();
+                self.render_data.node_children[node_element_index * 8 + 4] = empty_marker();
+                self.render_data.node_children[node_element_index * 8 + 5] = empty_marker();
+                self.render_data.node_children[node_element_index * 8 + 6] = empty_marker();
+                self.render_data.node_children[node_element_index * 8 + 7] = empty_marker();
+                #[cfg(debug_assertions)]
+                {
+                    if let BrickData::Solid(_) | BrickData::Empty = brick {
+                        // If no brick was added, the occupied bits should either be empty or full
+                        if let NodeChildrenArray::OccupancyBitmap(occupied_bits) =
+                            tree.node_children[node_key].content
+                        {
+                            debug_assert!(occupied_bits == 0 || occupied_bits == u64::MAX);
+                        }
+                    }
+                }
+            }
+            NodeContent::Leaf(bricks) => {
+                debug_assert!(
+                    matches!(
+                        tree.node_children[node_key].content,
+                        NodeChildrenArray::OccupancyBitmap(_)
+                    ),
+                    "Expected Leaf to have OccupancyBitmaps(_) instead of {:?}",
+                    tree.node_children[node_key].content
+                );
+                if try_add_children {
+                    for octant in 0..8 {
+                        let (brick_index, mut current_modified_nodes, mut current_modified_bricks) =
+                            self.add_brick(tree, node_key, octant);
+                        modified_bricks.push(brick_index as usize);
+                        modified_nodes.append(&mut current_modified_nodes);
+                        modified_bricks.append(&mut current_modified_bricks);
+
+                        self.render_data.node_children[node_element_index * 8 + octant] =
+                            brick_index;
+                        #[cfg(debug_assertions)]
+                        {
+                            if let BrickData::Solid(_) | BrickData::Empty = bricks[octant] {
+                                // If no brick was added, the relevant occupied bits should either be empty or full
+                                if let NodeChildrenArray::OccupancyBitmap(occupied_bits) =
+                                    tree.node_children[node_key].content
+                                {
+                                    debug_assert!(
+                                        0 == occupied_bits & BITMAP_MASK_FOR_OCTANT_LUT[octant]
+                                            || BITMAP_MASK_FOR_OCTANT_LUT[octant]
+                                                == occupied_bits
+                                                    & BITMAP_MASK_FOR_OCTANT_LUT[octant]
+                                    );
+                                }
+                            }
+                        }
+                    }
+                } else {
+                    for octant in 0..8 {
+                        self.render_data.node_children[node_element_index * 8 + octant] =
+                            empty_marker();
+                    }
+                }
+            }
+            NodeContent::Internal(_) => {
+                for octant in 0..8 {
+                    let child_key = tree.node_children[node_key][octant] as usize;
+                    if child_key != empty_marker() as usize {
+                        if try_add_children
+                            && !self.node_key_vs_meta_index.contains_left(&child_key)
+                        {
+                            // In case @try_add_children is true, no new node is added in case the cache is full,
+                            // so there will be no severed parents in this case
+                            self.add_node(tree, child_key, try_add_children);
+                        }
+
+                        self.render_data.node_children[node_element_index * 8 + octant as usize] =
+                            *self
+                                .node_key_vs_meta_index
+                                .get_by_left(&child_key)
+                                .unwrap_or(&(empty_marker() as usize))
+                                as u32;
+                    } else {
+                        self.render_data.node_children[node_element_index * 8 + octant as usize] =
+                            empty_marker();
+                    }
+                }
+            }
+            NodeContent::Nothing => {
+                for octant in 0..8 {
+                    self.render_data.node_children[node_element_index * 8 + octant as usize] =
+                        empty_marker();
+                }
+            }
+        }
+        Some((node_element_index, modified_nodes, modified_bricks))
+    }
+
+    //##############################################################################
+    //    █████████   ██████████   ██████████
+    //   ███░░░░░███ ░░███░░░░███ ░░███░░░░███
+    //  ░███    ░███  ░███   ░░███ ░███   ░░███
+    //  ░███████████  ░███    ░███ ░███    ░███
+    //  ░███░░░░░███  ░███    ░███ ░███    ░███
+    //  ░███    ░███  ░███    ███  ░███    ███
+    //  █████   █████ ██████████   ██████████
+    // ░░░░░   ░░░░░ ░░░░░░░░░░   ░░░░░░░░░░
+    //  ███████████  ███████████   █████   █████████  █████   ████
+    // ░░███░░░░░███░░███░░░░░███ ░░███   ███░░░░░███░░███   ███░
+    //  ░███    ░███ ░███    ░███  ░███  ███     ░░░  ░███  ███
+    //  ░██████████  ░██████████   ░███ ░███          ░███████
+    //  ░███░░░░░███ ░███░░░░░███  ░███ ░███          ░███░░███
+    //  ░███    ░███ ░███    ░███  ░███ ░░███     ███ ░███ ░░███
+    //  ███████████  █████   █████ █████ ░░█████████  █████ ░░████
+    // ░░░░░░░░░░░  ░░░░░   ░░░░░ ░░░░░   ░░░░░░░░░  ░░░░░   ░░░░
+    //##############################################################################
+    /// Provides the index of the first brick available to be overwritten, through the second chance algorithm
+    fn first_available_brick(&mut self) -> usize {
+        let mut brick_index;
+        loop {
+            brick_index = self.victim_brick;
+            if
+            // current brick is not owned or used
+            BrickOwnedBy::NotOwned == self.brick_ownership[brick_index]
+                || (0
+                    == (self.render_data.metadata[brick_index / 8]
+                        & Self::brick_used_mask(brick_index)))
+            {
+                // mark brick used
+                self.render_data.metadata[brick_index / 8] |= Self::brick_used_mask(brick_index);
+                break;
+            }
+
+            // mark current brick unused and step the iterator forward
+            self.render_data.metadata[brick_index / 8] &= !Self::brick_used_mask(brick_index);
+            self.victim_brick = (brick_index + 1) % (self.render_data.metadata.len() * 8);
+        }
+
+        brick_index
+    }
+
+    /// Loads a brick into the provided voxels vector and color palette
+    /// * `brick` - The brick to upload
+    /// * `tree` - The octree where the brick is found
+    /// * `returns` - the index where the brick is found and potentially a list of nodes and bricks modified during insertion
+    pub(crate) fn add_brick<T, const DIM: usize>(
+        &mut self,
+        tree: &Octree<T, DIM>,
+        node_key: usize,
+        target_octant: usize,
+    ) -> (u32, Vec<usize>, Vec<usize>)
+    where
+        T: Default + Clone + PartialEq + VoxelData + Send + Sync + 'static,
+    {
+        debug_assert_eq!(
+            self.render_data.voxels.len() % (DIM * DIM * DIM),
+            0,
+            "Expected Voxel buffer length({:?}) to be divisble by {:?}",
+            self.render_data.voxels.len(),
+            (DIM * DIM * DIM)
+        );
+
+        let brick = match tree.nodes.get(node_key) {
+            NodeContent::UniformLeaf(brick) => brick,
+            NodeContent::Leaf(bricks) => &bricks[target_octant],
+            NodeContent::Nothing | NodeContent::Internal(_) => {
+                panic!("Expected to add brick of Internal or empty node!")
+            }
+        };
+
+        match brick {
+            BrickData::Empty => (empty_marker(), Vec::new(), Vec::new()),
+            BrickData::Solid(voxel) => {
+                let albedo = voxel.albedo();
+                // The number of colors inserted into the palette is the size of the color palette map
+                let color_palette_size = self.map_to_color_index_in_palette.keys().len();
+                if let std::collections::hash_map::Entry::Vacant(e) =
+                    self.map_to_color_index_in_palette.entry(albedo)
+                {
+                    e.insert(color_palette_size);
+                    self.render_data.color_palette[color_palette_size] = Vec4::new(
+                        albedo.r as f32 / 255.,
+                        albedo.g as f32 / 255.,
+                        albedo.b as f32 / 255.,
+                        albedo.a as f32 / 255.,
+                    );
+                }
+                (
+                    self.map_to_color_index_in_palette[&albedo] as u32,
+                    Vec::new(),
+                    Vec::new(),
+                )
+            }
+            BrickData::Parted(brick) => {
+                if let Some(brick_index) = self
+                    .map_to_brick_maybe_owned_by_node
+                    .get(&(node_key, target_octant as u8))
+                {
+                    if self.brick_ownership[*brick_index] == BrickOwnedBy::NotOwned {
+                        self.brick_ownership[*brick_index] =
+                            BrickOwnedBy::Node(node_key as u32, target_octant as u8);
+                        return (*brick_index as u32, Vec::new(), Vec::new());
+                    } else {
+                        // remove from index if it is owned by another node already
+                        self.map_to_brick_maybe_owned_by_node
+                            .remove(&(node_key, target_octant as u8));
+                    }
+                }
+
+                let brick_index = self.first_available_brick();
+                let (modified_nodes, modified_bricks) =
+                    if let BrickOwnedBy::Node(key, octant) = self.brick_ownership[brick_index] {
+                        debug_assert!(
+                            self.node_key_vs_meta_index.contains_left(&(key as usize)),
+                            "Expected brick to be owned by a node used in cache"
+                        );
+
+                        self.erase_node_child(
+                            *self
+                                .node_key_vs_meta_index
+                                .get_by_left(&(key as usize))
+                                .unwrap(),
+                            octant as usize,
+                            tree,
+                        )
+                    } else {
+                        (Vec::new(), Vec::new())
+                    };
+
+                self.brick_ownership[brick_index as usize] =
+                    BrickOwnedBy::Node(node_key as u32, target_octant as u8);
+
+                for z in 0..DIM {
+                    for y in 0..DIM {
+                        for x in 0..DIM {
+                            // The number of colors inserted into the palette is the size of the color palette map
+                            let potential_new_albedo_index =
+                                self.map_to_color_index_in_palette.keys().len();
+                            let albedo = brick[x][y][z].albedo();
+                            let albedo_index = if let std::collections::hash_map::Entry::Vacant(e) =
+                                self.map_to_color_index_in_palette.entry(albedo)
+                            {
+                                e.insert(potential_new_albedo_index);
+                                self.render_data.color_palette[potential_new_albedo_index] =
+                                    Vec4::new(
+                                        albedo.r as f32 / 255.,
+                                        albedo.g as f32 / 255.,
+                                        albedo.b as f32 / 255.,
+                                        albedo.a as f32 / 255.,
+                                    );
+                                potential_new_albedo_index
+                            } else {
+                                self.map_to_color_index_in_palette[&albedo]
+                            };
+                            self.render_data.voxels[(brick_index * (DIM * DIM * DIM))
+                                + flat_projection(x, y, z, DIM)] = Voxelement {
+                                albedo_index: albedo_index as u32,
+                                content: brick[x][y][z].user_data(),
+                            };
+                        }
+                    }
+                }
+
+                (brick_index as u32, modified_nodes, modified_bricks)
+            }
+        }
+    }
+}
diff --git a/src/octree/raytracing/bevy/data.rs b/src/octree/raytracing/bevy/data.rs
index fa2fa5c..d2a7008 100644
--- a/src/octree/raytracing/bevy/data.rs
+++ b/src/octree/raytracing/bevy/data.rs
@@ -1,328 +1,565 @@
-use crate::{
-    object_pool::empty_marker,
-    octree::{
-        types::BrickData,
-        {
-            raytracing::bevy::types::{OctreeMetaData, ShocoVoxRenderData, Voxelement},
-            types::{NodeChildrenArray, NodeContent},
-            Albedo, Octree, V3c, VoxelData,
+use crate::object_pool::empty_marker;
+use crate::octree::{
+    raytracing::bevy::types::{
+        BrickOwnedBy, OctreeGPUDataHandler, OctreeGPUHost, OctreeGPUView, OctreeMetaData,
+        OctreeRenderData, OctreeSpyGlass, SvxRenderPipeline, SvxViewSet, VictimPointer, Viewport,
+        Voxelement,
+    },
+    BrickData, NodeContent, Octree, V3c, VoxelData,
+};
+use bevy::{
+    ecs::system::{Res, ResMut},
+    math::Vec4,
+    prelude::{Assets, Handle, Image},
+    render::{
+        render_asset::RenderAssetUsages,
+        render_resource::{
+            encase::{internal::WriteInto, StorageBuffer, UniformBuffer},
+            Buffer, Extent3d, ShaderSize, TextureDimension, TextureFormat, TextureUsages,
         },
+        renderer::{RenderDevice, RenderQueue},
     },
-    spatial::lut::BITMAP_MASK_FOR_OCTANT_LUT,
 };
-use bevy::math::Vec4;
-use std::collections::HashMap;
+use bimap::BiHashMap;
+use std::{
+    collections::{HashMap, HashSet},
+    sync::{Arc, Mutex},
+};
 
-impl<T, const DIM: usize> Octree<T, DIM>
+impl<T, const DIM: usize> OctreeGPUHost<T, DIM>
 where
-    T: Default + Clone + Copy + PartialEq + VoxelData,
+    T: Default + Clone + Copy + PartialEq + VoxelData + Send + Sync + 'static,
 {
-    /// Updates the meta element value to store that the corresponding node is a leaf node
-    fn meta_set_is_leaf(sized_node_meta: &mut u32, is_leaf: bool) {
-        *sized_node_meta =
-            (*sized_node_meta & 0xFFFFFFFB) | if is_leaf { 0x00000004 } else { 0x00000000 };
-    }
+    //##############################################################################
+    //     ███████      █████████  ███████████ ███████████   ██████████ ██████████
+    //   ███░░░░░███   ███░░░░░███░█░░░███░░░█░░███░░░░░███ ░░███░░░░░█░░███░░░░░█
+    //  ███     ░░███ ███     ░░░ ░   ░███  ░  ░███    ░███  ░███  █ ░  ░███  █ ░
+    // ░███      ░███░███             ░███     ░██████████   ░██████    ░██████
+    // ░███      ░███░███             ░███     ░███░░░░░███  ░███░░█    ░███░░█
+    // ░░███     ███ ░░███     ███    ░███     ░███    ░███  ░███ ░   █ ░███ ░   █
+    //  ░░░███████░   ░░█████████     █████    █████   █████ ██████████ ██████████
+    //    ░░░░░░░      ░░░░░░░░░     ░░░░░    ░░░░░   ░░░░░ ░░░░░░░░░░ ░░░░░░░░░░
+    //    █████████  ███████████  █████  █████
+    //   ███░░░░░███░░███░░░░░███░░███  ░░███
+    //  ███     ░░░  ░███    ░███ ░███   ░███
+    // ░███          ░██████████  ░███   ░███
+    // ░███    █████ ░███░░░░░░   ░███   ░███
+    // ░░███  ░░███  ░███         ░███   ░███
+    //  ░░█████████  █████        ░░████████
+    //   ░░░░░░░░░  ░░░░░          ░░░░░░░░
+    //  █████   █████ █████ ██████████ █████   ███   █████
+    // ░░███   ░░███ ░░███ ░░███░░░░░█░░███   ░███  ░░███
+    //  ░███    ░███  ░███  ░███  █ ░  ░███   ░███   ░███
+    //  ░███    ░███  ░███  ░██████    ░███   ░███   ░███
+    //  ░░███   ███   ░███  ░███░░█    ░░███  █████  ███
+    //   ░░░█████░    ░███  ░███ ░   █  ░░░█████░█████░
+    //     ░░███      █████ ██████████    ░░███ ░░███
+    //##############################################################################
 
-    /// Updates the meta element value to store that the corresponding node is a uniform leaf node
-    fn meta_set_is_uniform(sized_node_meta: &mut u32, is_uniform: bool) {
-        *sized_node_meta =
-            (*sized_node_meta & 0xFFFFFFF7) | if is_uniform { 0x00000008 } else { 0x00000000 };
-    }
+    /// Creates GPU compatible data renderable on the GPU from an octree
+    pub fn create_new_view(
+        &mut self,
+        svx_view_set: &mut SvxViewSet,
+        size: usize,
+        viewport: Viewport,
+        resolution: [u32; 2],
+        mut images: ResMut<Assets<Image>>,
+    ) -> Handle<Image> {
+        let mut gpu_data_handler = OctreeGPUDataHandler {
+            render_data: OctreeRenderData {
+                octree_meta: OctreeMetaData {
+                    octree_size: self.tree.octree_size,
+                    voxel_brick_dim: DIM as u32,
+                    ambient_light_color: V3c::new(1., 1., 1.),
+                    ambient_light_position: V3c::new(
+                        self.tree.octree_size as f32,
+                        self.tree.octree_size as f32,
+                        self.tree.octree_size as f32,
+                    ),
+                },
+                metadata: vec![0; size],
+                node_ocbits: vec![0; size * 2],
+                node_children: vec![empty_marker(); size * 8],
+                color_palette: vec![Vec4::ZERO; u16::MAX as usize],
+                voxels: vec![
+                    Voxelement {
+                        albedo_index: 0,
+                        content: 0
+                    };
+                    size * 8 * (DIM * DIM * DIM)
+                ],
+            },
+            victim_node: VictimPointer::new(size),
+            victim_brick: 0,
+            map_to_color_index_in_palette: HashMap::new(),
+            map_to_brick_maybe_owned_by_node: HashMap::new(),
+            node_key_vs_meta_index: BiHashMap::new(),
+            brick_ownership: vec![BrickOwnedBy::NotOwned; size * 8],
+            uploaded_color_palette_size: 0,
+        };
 
-    /// Updates the meta element value to store the brick structure of the given leaf node.
-    /// Does not erase anything in @sized_node_meta, it's expected to be cleared before
-    /// the first use of this function
-    /// for the given brick octant
-    /// * `sized_node_meta` - the bytes to update
-    /// * `brick` - the brick to describe into the bytes
-    /// * `brick_octant` - the octant to update in the bytes
-    fn meta_add_leaf_brick_structure(
-        sized_node_meta: &mut u32,
-        brick: &BrickData<T, DIM>,
-        brick_octant: usize,
-    ) {
-        match brick {
-            BrickData::Empty => {} // Child structure properties already set to NIL
-            BrickData::Solid(_voxel) => {
-                // set child Occupied bits, child Structure bits already set to NIL
-                *sized_node_meta |= 0x01 << (8 + brick_octant);
-            }
-            BrickData::Parted(_brick) => {
-                // set child Occupied bits
-                *sized_node_meta |= 0x01 << (8 + brick_octant);
+        gpu_data_handler.add_node(&self.tree, Octree::<T, DIM>::ROOT_NODE_KEY as usize, true);
 
-                // set child Structure bits
-                *sized_node_meta |= 0x01 << (16 + brick_octant);
-            }
-        };
+        let mut output_texture = Image::new_fill(
+            Extent3d {
+                width: resolution[0],
+                height: resolution[1],
+                depth_or_array_layers: 1,
+            },
+            TextureDimension::D2,
+            &[0, 0, 0, 255],
+            TextureFormat::Rgba8Unorm,
+            RenderAssetUsages::RENDER_WORLD,
+        );
+        output_texture.texture_descriptor.usage = TextureUsages::COPY_DST
+            | TextureUsages::STORAGE_BINDING
+            | TextureUsages::TEXTURE_BINDING;
+        let output_texture = images.add(output_texture);
+
+        svx_view_set.views.push(Arc::new(Mutex::new(OctreeGPUView {
+            data_handler: gpu_data_handler,
+            spyglass: OctreeSpyGlass {
+                node_requests: vec![empty_marker(); 4],
+                output_texture: output_texture.clone(),
+                viewport: viewport,
+            },
+        })));
+        output_texture
     }
+}
 
-    /// Updates the given meta element value to store the leaf structure of the given node
-    /// the given NodeContent reference is expected to be a leaf node
-    fn meta_set_leaf_structure(sized_node_meta: &mut u32, leaf: &NodeContent<T, DIM>) {
-        match leaf {
-            NodeContent::UniformLeaf(brick) => {
-                Self::meta_set_is_leaf(sized_node_meta, true);
-                Self::meta_set_is_uniform(sized_node_meta, true);
-                Self::meta_add_leaf_brick_structure(sized_node_meta, brick, 0);
-            }
-            NodeContent::Leaf(bricks) => {
-                Self::meta_set_is_leaf(sized_node_meta, true);
-                Self::meta_set_is_uniform(sized_node_meta, false);
-                for octant in 0..8 {
-                    Self::meta_add_leaf_brick_structure(sized_node_meta, &bricks[octant], octant);
+/// Handles data sync between Bevy main(CPU) world and rendering world
+pub(crate) fn sync_with_main_world(// tree_view: Option<ResMut<OctreeGPUView>>,
+    // mut world: ResMut<bevy::render::MainWorld>,
+) {
+    // This function is unused because ExtractResource plugin is handling the sync
+    // However, it is only one way: MainWorld --> RenderWorld
+    // Any modification here is overwritten by the plugin if it is active,
+    // in order to enable data flow in the opposite direction, extractresource plugin
+    // needs to be disabled, and the sync logic (both ways) needs to be implemented here
+    // refer to: https://www.reddit.com/r/bevy/comments/1ay50ee/copy_from_render_world_to_main_world/
+}
+
+//##############################################################################
+//    █████████  ███████████  █████  █████
+//   ███░░░░░███░░███░░░░░███░░███  ░░███
+//  ███     ░░░  ░███    ░███ ░███   ░███
+// ░███          ░██████████  ░███   ░███
+// ░███    █████ ░███░░░░░░   ░███   ░███
+// ░░███  ░░███  ░███         ░███   ░███
+//  ░░█████████  █████        ░░████████
+//   ░░░░░░░░░  ░░░░░          ░░░░░░░░
+//  ███████████   ██████████   █████████   ██████████
+// ░░███░░░░░███ ░░███░░░░░█  ███░░░░░███ ░░███░░░░███
+//  ░███    ░███  ░███  █ ░  ░███    ░███  ░███   ░░███
+//  ░██████████   ░██████    ░███████████  ░███    ░███
+//  ░███░░░░░███  ░███░░█    ░███░░░░░███  ░███    ░███
+//  ░███    ░███  ░███ ░   █ ░███    ░███  ░███    ███
+//  █████   █████ ██████████ █████   █████ ██████████
+// ░░░░░   ░░░░░ ░░░░░░░░░░ ░░░░░   ░░░░░ ░░░░░░░░░░
+//##############################################################################
+/// Handles data reads from GPU every loop, mainly data requests and usaage updates.
+/// Based on https://docs.rs/bevy/latest/src/gpu_readback/gpu_readback.rs.html
+pub(crate) fn handle_gpu_readback<T, const DIM: usize>(
+    render_device: Res<RenderDevice>,
+    svx_view_set: ResMut<SvxViewSet>,
+    mut svx_pipeline: Option<ResMut<SvxRenderPipeline>>,
+) where
+    T: Default + Clone + PartialEq + VoxelData + Send + Sync + 'static,
+{
+    if let Some(ref mut pipeline) = svx_pipeline {
+        let resources = pipeline.resources.as_ref().unwrap();
+
+        let node_requests_buffer_slice = resources.readable_node_requests_buffer.slice(..);
+        let (s, node_requests_recv) = crossbeam::channel::unbounded::<()>();
+        node_requests_buffer_slice.map_async(
+            bevy::render::render_resource::MapMode::Read,
+            move |d| match d {
+                Ok(_) => s.send(()).expect("Failed to send map update"),
+                Err(err) => panic!("Couldn't map debug interface buffer!: {err}"),
+            },
+        );
+
+        render_device
+            .poll(bevy::render::render_resource::Maintain::wait())
+            .panic_on_timeout();
+
+        let mut view = svx_view_set.views[0].lock().unwrap();
+        node_requests_recv
+            .recv()
+            .expect("Failed to receive the map_async message");
+        {
+            let buffer_view = node_requests_buffer_slice.get_mapped_range();
+            view.spyglass.node_requests = buffer_view
+                .chunks(std::mem::size_of::<u32>())
+                .map(|chunk| u32::from_ne_bytes(chunk.try_into().expect("should be a u32")))
+                .collect::<Vec<u32>>();
+        }
+        resources.readable_node_requests_buffer.unmap();
+
+        if {
+            let mut is_metadata_required_this_loop = false;
+            for node_request in &view.spyglass.node_requests {
+                if *node_request != empty_marker() {
+                    is_metadata_required_this_loop = true;
+                    break;
                 }
             }
-            NodeContent::Internal(_) | NodeContent::Nothing => {
-                panic!("Expected node content to be of a leaf");
+            is_metadata_required_this_loop
+        } {
+            let metadata_buffer_slice = resources.readable_metadata_buffer.slice(..);
+            let (s, metadata_recv) = crossbeam::channel::unbounded::<()>();
+            metadata_buffer_slice.map_async(
+                bevy::render::render_resource::MapMode::Read,
+                move |d| match d {
+                    Ok(_) => s.send(()).expect("Failed to send map update"),
+                    Err(err) => panic!("Couldn't map debug interface buffer!: {err}"),
+                },
+            );
+
+            render_device
+                .poll(bevy::render::render_resource::Maintain::wait())
+                .panic_on_timeout();
+            metadata_recv
+                .recv()
+                .expect("Failed to receive the map_async message");
+            {
+                let buffer_view = metadata_buffer_slice.get_mapped_range();
+                view.data_handler.render_data.metadata = buffer_view
+                    .chunks(std::mem::size_of::<u32>())
+                    .map(|chunk| u32::from_ne_bytes(chunk.try_into().expect("should be a u32")))
+                    .collect::<Vec<u32>>();
             }
+            resources.readable_metadata_buffer.unmap();
         }
     }
+}
 
-    /// Creates the descriptor bytes for the given node
-    fn create_node_properties(node: &NodeContent<T, DIM>) -> u32 {
-        let mut meta = 0;
-        match node {
-            NodeContent::Leaf(_) | NodeContent::UniformLeaf(_) => {
-                Self::meta_set_is_leaf(&mut meta, true);
-                Self::meta_set_leaf_structure(&mut meta, node);
-            }
-            NodeContent::Internal(_) | NodeContent::Nothing => {
-                Self::meta_set_is_leaf(&mut meta, false);
-            }
-        };
-        meta
-    }
+//##############################################################################
+//    █████████  ███████████  █████  █████
+//   ███░░░░░███░░███░░░░░███░░███  ░░███
+//  ███     ░░░  ░███    ░███ ░███   ░███
+// ░███          ░██████████  ░███   ░███
+// ░███    █████ ░███░░░░░░   ░███   ░███
+// ░░███  ░░███  ░███         ░███   ░███
+//  ░░█████████  █████        ░░████████
+//   ░░░░░░░░░  ░░░░░          ░░░░░░░░
 
-    /// Loads a brick into the provided voxels vector and color palette
-    /// * `brick` - The brick to upload
-    /// * `voxels` - The destination buffer
-    /// * `color_palette` - The used color palette
-    /// * `map_to_color_index_in_palette` - Indexing helper for the color palette
-    /// * `returns` - the identifier to set in @SizedNode and true if a new brick was aded to the voxels vector
-    fn add_brick_to_vec(
-        brick: &BrickData<T, DIM>,
-        voxels: &mut Vec<Voxelement>,
-        color_palette: &mut Vec<Vec4>,
-        map_to_color_index_in_palette: &mut HashMap<Albedo, usize>,
-    ) -> (u32, bool) {
-        match brick {
-            BrickData::Empty => (empty_marker(), false),
-            BrickData::Solid(voxel) => {
-                let albedo = voxel.albedo();
-                if let std::collections::hash_map::Entry::Vacant(e) =
-                    map_to_color_index_in_palette.entry(albedo)
-                {
-                    e.insert(color_palette.len());
-                    color_palette.push(Vec4::new(
-                        albedo.r as f32 / 255.,
-                        albedo.g as f32 / 255.,
-                        albedo.b as f32 / 255.,
-                        albedo.a as f32 / 255.,
-                    ));
-                }
-                (map_to_color_index_in_palette[&albedo] as u32, false)
-            }
-            BrickData::Parted(brick) => {
-                voxels.reserve(DIM * DIM * DIM);
-                let brick_index = voxels.len() / (DIM * DIM * DIM);
-                debug_assert_eq!(
-                    voxels.len() % (DIM * DIM * DIM),
-                    0,
-                    "Expected Voxel buffer length({:?}) to be divisble by {:?}",
-                    voxels.len(),
-                    (DIM * DIM * DIM)
-                );
-                for z in 0..DIM {
-                    for y in 0..DIM {
-                        for x in 0..DIM {
-                            let albedo = brick[x][y][z].albedo();
-                            if let std::collections::hash_map::Entry::Vacant(e) =
-                                map_to_color_index_in_palette.entry(albedo)
-                            {
-                                e.insert(color_palette.len());
-                                color_palette.push(Vec4::new(
-                                    albedo.r as f32 / 255.,
-                                    albedo.g as f32 / 255.,
-                                    albedo.b as f32 / 255.,
-                                    albedo.a as f32 / 255.,
-                                ));
-                            }
-                            let albedo_index = map_to_color_index_in_palette[&albedo];
+//  █████   ███   █████ ███████████   █████ ███████████ ██████████
+// ░░███   ░███  ░░███ ░░███░░░░░███ ░░███ ░█░░░███░░░█░░███░░░░░█
+//  ░███   ░███   ░███  ░███    ░███  ░███ ░   ░███  ░  ░███  █ ░
+//  ░███   ░███   ░███  ░██████████   ░███     ░███     ░██████
+//  ░░███  █████  ███   ░███░░░░░███  ░███     ░███     ░███░░█
+//   ░░░█████░█████░    ░███    ░███  ░███     ░███     ░███ ░   █
+//     ░░███ ░░███      █████   █████ █████    █████    ██████████
+//      ░░░   ░░░      ░░░░░   ░░░░░ ░░░░░    ░░░░░    ░░░░░░░░░░
+//##############################################################################
 
-                            voxels.push(Voxelement {
-                                albedo_index: albedo_index as u32,
-                                content: brick[x][y][z].user_data(),
-                            });
-                        }
-                    }
-                }
-                (brick_index as u32, true)
-            }
+/// Converts the given array to `&[u8]` on the given range,
+/// and schedules it to be written to the given buffer in the GPU
+fn write_range_to_buffer<U>(
+    array: &Vec<U>,
+    range: std::ops::Range<usize>,
+    buffer: &Buffer,
+    render_queue: &RenderQueue,
+) where
+    U: Send + Sync + 'static + ShaderSize + WriteInto,
+{
+    if !range.is_empty() {
+        let element_size = std::mem::size_of_val(&array[0]);
+        let byte_offset = (range.start * element_size) as u64;
+        let slice = array.get(range.clone()).expect(
+            &format!(
+                "Expected range {:?} to be in bounds of {:?}",
+                range,
+                array.len(),
+            )
+            .to_owned(),
+        );
+        unsafe {
+            render_queue.write_buffer(buffer, byte_offset, &slice.align_to::<u8>().1);
         }
     }
+}
 
-    /// Creates GPU compatible data renderable on the GPU from an octree
-    pub fn create_bevy_view(&self) -> ShocoVoxRenderData {
-        let meta = OctreeMetaData {
-            octree_size: self.octree_size,
-            voxel_brick_dim: DIM as u32,
-            ambient_light_color: V3c::new(1., 1., 1.),
-            ambient_light_position: V3c::new(
-                self.octree_size as f32,
-                self.octree_size as f32,
-                self.octree_size as f32,
-            ),
+/// Handles Data Streaming to the GPU based on incoming requests from the view(s)
+pub(crate) fn write_to_gpu<T, const DIM: usize>(
+    tree_gpu_host: Option<ResMut<OctreeGPUHost<T, DIM>>>,
+    svx_pipeline: Option<ResMut<SvxRenderPipeline>>,
+    svx_view_set: ResMut<SvxViewSet>,
+) where
+    T: Default + Clone + Copy + PartialEq + VoxelData + Send + Sync + 'static,
+{
+    if let (Some(pipeline), Some(tree_host)) = (svx_pipeline, tree_gpu_host) {
+        let render_queue = &pipeline.render_queue;
+        let resources = if let Some(resources) = &pipeline.resources {
+            resources
+        } else {
+            // No resources available yet, can't write to them
+            return;
         };
 
-        let mut nodes = Vec::new();
-        let mut children_buffer = Vec::with_capacity(self.nodes.len() * 8);
-        let mut voxels = Vec::new();
-        let mut node_occupied_bits = Vec::new();
-        let mut color_palette = Vec::new();
-
-        // Build up Nodes
-        let mut map_to_node_index_in_nodes_buffer = HashMap::new();
-        for i in 0..self.nodes.len() {
-            if self.nodes.key_is_valid(i) {
-                map_to_node_index_in_nodes_buffer.insert(i, nodes.len());
-                nodes.push(Self::create_node_properties(self.nodes.get(i)));
-            }
-        }
+        let mut view = svx_view_set.views[0].lock().unwrap();
 
-        // Build up voxel content
-        let mut map_to_color_index_in_palette = HashMap::new();
-        for i in 0..self.nodes.len() {
-            if !self.nodes.key_is_valid(i) {
-                continue;
-            }
-            let occupied_bits = self.stored_occupied_bits(i);
-            node_occupied_bits.extend_from_slice(&[
-                (occupied_bits & 0x00000000FFFFFFFF) as u32,
-                ((occupied_bits & 0xFFFFFFFF00000000) >> 32) as u32,
-            ]);
-            match self.nodes.get(i) {
-                NodeContent::UniformLeaf(brick) => {
-                    debug_assert!(
-                        matches!(
-                            self.node_children[i].content,
-                            NodeChildrenArray::OccupancyBitmap(_)
-                        ),
-                        "Expected Uniform leaf to have OccupancyBitmap(_) instead of {:?}",
-                        self.node_children[i].content
-                    );
-
-                    let (brick_index, brick_added) = Self::add_brick_to_vec(
-                        brick,
-                        &mut voxels,
-                        &mut color_palette,
-                        &mut map_to_color_index_in_palette,
-                    );
-
-                    children_buffer.extend_from_slice(&[
-                        brick_index,
-                        empty_marker(),
-                        empty_marker(),
-                        empty_marker(),
-                        empty_marker(),
-                        empty_marker(),
-                        empty_marker(),
-                        empty_marker(),
-                    ]);
-                    #[cfg(debug_assertions)]
-                    {
-                        if !brick_added {
-                            // If no brick was added, the occupied bits should either be empty or full
-                            if let NodeChildrenArray::OccupancyBitmap(occupied_bits) =
-                                self.node_children[i].content
-                            {
-                                debug_assert!(occupied_bits == 0 || occupied_bits == u64::MAX);
-                            }
-                        }
-                    }
+        // Data updates for spyglass viewport
+        let mut buffer = UniformBuffer::new(Vec::<u8>::new());
+        buffer.write(&view.spyglass.viewport).unwrap();
+        render_queue.write_buffer(&resources.viewport_buffer, 0, &buffer.into_inner());
+
+        // Handle node requests, update cache
+        let tree = &tree_host.tree;
+        {
+            let mut meta_updated = std::ops::Range {
+                start: view.data_handler.render_data.metadata.len(),
+                end: 0,
+            };
+            let mut ocbits_updated = std::ops::Range {
+                start: view.data_handler.render_data.node_ocbits.len(),
+                end: 0,
+            };
+            let mut node_children_updated = std::ops::Range {
+                start: view.data_handler.render_data.node_children.len(),
+                end: 0,
+            };
+            let mut voxels_updated = std::ops::Range {
+                start: view.data_handler.render_data.voxels.len(),
+                end: 0,
+            };
+            let mut node_requests = view.spyglass.node_requests.clone();
+            let mut modified_nodes = HashSet::<usize>::new();
+            let mut modified_bricks = HashSet::<usize>::new();
+            let victim_node_loop_count = view.data_handler.victim_node.get_loop_count();
+            for node_request in &mut node_requests {
+                if *node_request == empty_marker() {
+                    continue;
                 }
-                NodeContent::Leaf(bricks) => {
-                    debug_assert!(
-                        matches!(
-                            self.node_children[i].content,
-                            NodeChildrenArray::OccupancyBitmap(_)
-                        ),
-                        "Expected Leaf to have OccupancyBitmaps(_) instead of {:?}",
-                        self.node_children[i].content
-                    );
-
-                    let mut children = vec![empty_marker(); 8];
-                    for octant in 0..8 {
-                        let (brick_index, brick_added) = Self::add_brick_to_vec(
-                            &bricks[octant],
-                            &mut voxels,
-                            &mut color_palette,
-                            &mut map_to_color_index_in_palette,
+                let requested_parent_meta_index = (*node_request & 0x00FFFFFF) as usize;
+                let requested_child_octant = (*node_request & 0xFF000000) >> 24;
+
+                if modified_nodes.contains(&requested_parent_meta_index) {
+                    // Do not accept a request if the requester meta is already overwritten
+                    continue;
+                }
+
+                debug_assert!(view
+                    .data_handler
+                    .node_key_vs_meta_index
+                    .contains_right(&requested_parent_meta_index));
+                let requested_parent_node_key = view
+                    .data_handler
+                    .node_key_vs_meta_index
+                    .get_by_right(&requested_parent_meta_index)
+                    .unwrap()
+                    .clone();
+
+                debug_assert!(
+                    tree.nodes.key_is_valid(requested_parent_node_key),
+                    "Expected parent node({:?}) to be valid in GPU request.",
+                    requested_parent_node_key
+                );
+
+                modified_nodes.insert(requested_parent_meta_index);
+                match tree.nodes.get(requested_parent_node_key) {
+                    NodeContent::Nothing => {} // parent is empty, nothing to do
+                    NodeContent::Internal(_) => {
+                        let requested_child_node_key = tree_host.tree.node_children
+                            [requested_parent_node_key][requested_child_octant]
+                            as usize;
+                        debug_assert!(
+                            tree.nodes.key_is_valid(requested_child_node_key),
+                            "Expected key({:?}, child of node[{:?}][{:?}] in meta[{:?}]) to be valid in GPU request.",
+                            requested_child_node_key, requested_parent_node_key, requested_child_octant, requested_parent_meta_index
                         );
+                        let child_index = if !view
+                            .data_handler
+                            .node_key_vs_meta_index
+                            .contains_left(&requested_child_node_key)
+                        {
+                            let (child_index, currently_modified_nodes, currently_modified_bricks) =
+                                view.data_handler
+                                .add_node(&tree, requested_child_node_key, false)
+                                .expect("Expected to succeed adding a node into the GPU cache through data_handler");
+                            modified_nodes.extend(currently_modified_nodes);
+                            modified_bricks.extend(currently_modified_bricks);
+
+                            child_index
+                        } else {
+                            *view
+                                .data_handler
+                                .node_key_vs_meta_index
+                                .get_by_left(&requested_child_node_key)
+                                .unwrap()
+                        };
 
-                        children[octant] = brick_index;
-                        #[cfg(debug_assertions)]
+                        // Update connection to parent
+                        view.data_handler.render_data.node_children
+                            [requested_parent_meta_index * 8 + requested_child_octant as usize] =
+                            child_index as u32;
+
+                        debug_assert!(
+                            view.data_handler
+                                .node_key_vs_meta_index
+                                .contains_right(&requested_parent_meta_index),
+                            "Requester parent erased while adding its child node to meta"
+                        );
+
+                        ocbits_updated.start = ocbits_updated.start.min(child_index * 2);
+                        ocbits_updated.end = ocbits_updated.end.max(child_index * 2 + 2);
+                    }
+                    NodeContent::UniformLeaf(brick) => {
+                        // Only upload brick if it's not already available
+                        if matches!(brick, BrickData::Parted(_) | BrickData::Solid(_))
+                            && view.data_handler.render_data.node_children
+                                [requested_parent_meta_index * 8]
+                                == empty_marker()
                         {
-                            if !brick_added {
-                                // If no brick was added, the relevant occupied bits should either be empty or full
-                                if let NodeChildrenArray::OccupancyBitmap(occupied_bits) =
-                                    self.node_children[i].content
-                                {
-                                    debug_assert!(
-                                        0 == occupied_bits & BITMAP_MASK_FOR_OCTANT_LUT[octant]
-                                            || BITMAP_MASK_FOR_OCTANT_LUT[octant]
-                                                == occupied_bits
-                                                    & BITMAP_MASK_FOR_OCTANT_LUT[octant]
-                                    );
-                                }
+                            let (brick_index, currently_modified_nodes, currently_modified_bricks) =
+                                view.data_handler
+                                    .add_brick(&tree, requested_parent_node_key, 0);
+                            view.data_handler.render_data.node_children
+                                [requested_parent_meta_index * 8] = brick_index;
+
+                            modified_nodes.extend(currently_modified_nodes);
+                            modified_bricks.extend(currently_modified_bricks);
+
+                            if let BrickData::Parted(_) = brick {
+                                voxels_updated.start = voxels_updated
+                                    .start
+                                    .min(brick_index as usize * (DIM * DIM * DIM));
+                                voxels_updated.end = voxels_updated.end.max(
+                                    brick_index as usize * (DIM * DIM * DIM) + (DIM * DIM * DIM),
+                                );
                             }
                         }
                     }
-                    children_buffer.extend_from_slice(&children);
-                }
-                NodeContent::Internal(_) => {
-                    for c in 0..8 {
-                        let child_index = &self.node_children[i][c];
-                        if *child_index != empty_marker() {
-                            debug_assert!(map_to_node_index_in_nodes_buffer
-                                .contains_key(&(*child_index as usize)));
-                            children_buffer.push(
-                                map_to_node_index_in_nodes_buffer[&(*child_index as usize)] as u32,
-                            );
-                        } else {
-                            children_buffer.push(*child_index);
+                    NodeContent::Leaf(bricks) => {
+                        // Only upload brick if it's not already available
+                        if matches!(
+                            bricks[requested_child_octant as usize],
+                            BrickData::Parted(_) | BrickData::Solid(_)
+                        ) && view.data_handler.render_data.node_children
+                            [requested_parent_meta_index * 8 + requested_child_octant as usize]
+                            == empty_marker()
+                        {
+                            let (brick_index, currently_modified_nodes, currently_modified_bricks) =
+                                view.data_handler.add_brick(
+                                    &tree,
+                                    requested_parent_node_key,
+                                    requested_child_octant as usize,
+                                );
+                            view.data_handler.render_data.node_children[requested_parent_meta_index
+                                * 8
+                                + requested_child_octant as usize] = brick_index;
+
+                            modified_nodes.extend(currently_modified_nodes);
+                            modified_bricks.extend(currently_modified_bricks);
+
+                            if let BrickData::Parted(_) = bricks[requested_child_octant as usize] {
+                                voxels_updated.start = voxels_updated
+                                    .start
+                                    .min(brick_index as usize * (DIM * DIM * DIM));
+                                voxels_updated.end = voxels_updated.end.max(
+                                    brick_index as usize * (DIM * DIM * DIM) + (DIM * DIM * DIM),
+                                );
+                            }
                         }
                     }
                 }
-                NodeContent::Nothing => {
-                    children_buffer.extend_from_slice(&[empty_marker(); 8]);
+
+                if victim_node_loop_count != view.data_handler.victim_node.get_loop_count() {
+                    break;
                 }
             }
-        }
 
-        debug_assert_eq!(
-            nodes.len() * 2,
-            node_occupied_bits.len(),
-            "Node occupancy bitmaps length({:?}) should match node count({:?})!",
-            node_occupied_bits.len(),
-            nodes.len(),
-        );
+            debug_assert!(
+                // Either all node requests are empty
+                node_requests
+                    .iter()
+                    .filter(|&v| *v == empty_marker())
+                    .count()
+                    == node_requests.len()
+                    // Or some ndoes were updated this loop
+                    || 0 < modified_nodes.len()
+                    // Or the distance traveled by the victim pointer this loop is small enough
+                    || (view.data_handler.victim_node.len() as f32 * 0.5) as usize
+                        > (victim_node_loop_count - view.data_handler.victim_node.get_loop_count()),
+                "Couldn't process a single request because size of the buffer is too small."
+            );
 
-        debug_assert_eq!(
-            nodes.len() * 8,
-            children_buffer.len(),
-            "Node count({:?}) should match length of children buffer({:?})!",
-            nodes.len(),
-            children_buffer.len()
-        );
+            for node_request in &mut node_requests {
+                *node_request = empty_marker();
+            }
+
+            // Set updated buffers range based on modified nodes and bricks
+            for modified_node_index in &modified_nodes {
+                meta_updated.start = meta_updated.start.min(*modified_node_index);
+                meta_updated.end = meta_updated.end.max(modified_node_index + 1);
+                node_children_updated.start =
+                    node_children_updated.start.min(modified_node_index * 8);
+                node_children_updated.end =
+                    node_children_updated.end.max(modified_node_index * 8 + 8);
+            }
+
+            for modified_brick_index in &modified_bricks {
+                meta_updated.start = meta_updated.start.min(modified_brick_index / 8);
+                meta_updated.end = meta_updated.end.max(modified_brick_index / 8 + 1);
+            }
+
+            // write back updated data
+            let host_color_count = view.data_handler.map_to_color_index_in_palette.keys().len();
+            let color_palette_size_diff =
+                host_color_count - view.data_handler.uploaded_color_palette_size;
+            let resources = &pipeline.resources.as_ref().unwrap();
+
+            debug_assert!(
+                host_color_count >= view.data_handler.uploaded_color_palette_size,
+                "Expected host color palette({:?}), to be larger, than colors stored on the GPU({:?})",
+                host_color_count, view.data_handler.uploaded_color_palette_size
+            );
+            view.data_handler.uploaded_color_palette_size =
+                view.data_handler.map_to_color_index_in_palette.keys().len();
+
+            // Node requests
+            let mut buffer = StorageBuffer::new(Vec::<u8>::new());
+            buffer.write(&node_requests).unwrap();
+            render_queue.write_buffer(&resources.node_requests_buffer, 0, &buffer.into_inner());
+
+            // Color palette
+            if 0 < color_palette_size_diff {
+                // Upload color palette delta to GPU
+                write_range_to_buffer(
+                    &view.data_handler.render_data.color_palette,
+                    (host_color_count - color_palette_size_diff)..(host_color_count),
+                    &resources.color_palette_buffer,
+                    &render_queue,
+                );
+            }
 
-        ShocoVoxRenderData {
-            meta,
-            nodes,
-            children_buffer,
-            voxels,
-            node_occupied_bits,
-            color_palette,
+            // Render data
+            write_range_to_buffer(
+                &view.data_handler.render_data.metadata,
+                meta_updated,
+                &resources.metadata_buffer,
+                &render_queue,
+            );
+            write_range_to_buffer(
+                &view.data_handler.render_data.node_children,
+                node_children_updated,
+                &resources.node_children_buffer,
+                &render_queue,
+            );
+            write_range_to_buffer(
+                &view.data_handler.render_data.node_ocbits,
+                ocbits_updated,
+                &resources.node_ocbits_buffer,
+                &render_queue,
+            );
+            write_range_to_buffer(
+                &view.data_handler.render_data.voxels,
+                voxels_updated,
+                &resources.voxels_buffer,
+                &render_queue,
+            );
         }
     }
 }
diff --git a/src/octree/raytracing/bevy/mod.rs b/src/octree/raytracing/bevy/mod.rs
index cbdb072..123e9ee 100644
--- a/src/octree/raytracing/bevy/mod.rs
+++ b/src/octree/raytracing/bevy/mod.rs
@@ -1,204 +1,74 @@
+mod cache;
 mod data;
+mod pipeline;
 pub mod types;
 
 pub use crate::octree::raytracing::bevy::types::{
-    ShocoVoxRenderPlugin, ShocoVoxViewingGlass, Viewport,
+    OctreeGPUHost, OctreeGPUView, OctreeSpyGlass, RenderBevyPlugin, SvxViewSet, Viewport,
 };
 
-use crate::octree::raytracing::bevy::types::{
-    ShocoVoxLabel, ShocoVoxRenderData, ShocoVoxRenderNode, ShocoVoxRenderPipeline,
+use crate::octree::{
+    raytracing::bevy::{
+        data::{handle_gpu_readback, sync_with_main_world, write_to_gpu},
+        pipeline::prepare_bind_groups,
+        types::{SvxLabel, SvxRenderNode, SvxRenderPipeline},
+    },
+    VoxelData,
 };
 
 use bevy::{
     app::{App, Plugin},
-    asset::{AssetServer, Assets, Handle},
-    ecs::system::{Res, ResMut},
-    ecs::world::{FromWorld, World},
-    prelude::IntoSystemConfigs,
+    prelude::{ExtractSchedule, IntoSystemConfigs},
     render::{
-        extract_resource::ExtractResourcePlugin,
-        prelude::Image,
-        render_asset::{RenderAssetUsages, RenderAssets},
-        render_graph,
-        render_graph::RenderGraph,
-        render_resource::{
-            AsBindGroup, CachedPipelineState, ComputePassDescriptor, ComputePipelineDescriptor,
-            Extent3d, PipelineCache, TextureDimension, TextureFormat, TextureUsages,
-        },
-        renderer::{RenderContext, RenderDevice},
-        texture::{FallbackImage, GpuImage},
-        Render, RenderApp, RenderSet,
+        extract_resource::ExtractResourcePlugin, render_graph::RenderGraph, Render, RenderApp,
+        RenderSet,
     },
 };
 
-use std::borrow::Cow;
-
-pub fn create_viewing_glass(
-    viewport: &Viewport,
-    resolution: [u32; 2],
-    images: ResMut<Assets<Image>>,
-) -> 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::<RenderDevice>();
-        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::<AssetServer>()
-            .load("shaders/viewport_render.wgsl");
-        let pipeline_cache = world.resource::<PipelineCache>();
-        let update_pipeline = pipeline_cache.queue_compute_pipeline(ComputePipelineDescriptor {
-            label: None,
-            layout: vec![
-                viewing_glass_bind_group_layout.clone(),
-                render_data_bind_group_layout.clone(),
-            ],
-            push_constant_ranges: Vec::new(),
-            shader,
-            shader_defs: vec![],
-            entry_point: Cow::from("update"),
-        });
-
-        ShocoVoxRenderPipeline {
-            update_tree: true,
-            viewing_glass_bind_group_layout,
-            render_data_bind_group_layout,
-            update_pipeline,
-            viewing_glass_bind_group: None,
-            tree_bind_group: None,
+impl<T, const DIM: usize> RenderBevyPlugin<T, DIM>
+where
+    T: Default + Clone + PartialEq + VoxelData + Send + Sync + 'static,
+{
+    pub fn new(resolution: [u32; 2]) -> Self {
+        RenderBevyPlugin {
+            dummy: std::marker::PhantomData,
+            resolution,
         }
     }
 }
 
-fn prepare_bind_groups(
-    gpu_images: Res<RenderAssets<GpuImage>>,
-    fallback_image: Res<FallbackImage>,
-    render_device: Res<RenderDevice>,
-    mut pipeline: ResMut<ShocoVoxRenderPipeline>,
-    octree_viewing_glass: Res<ShocoVoxViewingGlass>,
-    render_data: Res<ShocoVoxRenderData>,
-) {
-    let bind_group = octree_viewing_glass
-        .as_bind_group(
-            &pipeline.viewing_glass_bind_group_layout,
-            &render_device,
-            &gpu_images,
-            &fallback_image,
-        )
-        .ok()
-        .unwrap();
-    pipeline.viewing_glass_bind_group = Some(bind_group.bind_group);
-
-    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(
-    resolution: [u32; 2],
-    mut images: ResMut<Assets<Image>>,
-) -> Handle<Image> {
-    let mut output_texture = Image::new_fill(
-        Extent3d {
-            width: resolution[0],
-            height: resolution[1],
-            depth_or_array_layers: 1,
-        },
-        TextureDimension::D2,
-        &[0, 0, 0, 255],
-        TextureFormat::Rgba8Unorm,
-        RenderAssetUsages::RENDER_WORLD,
-    );
-    output_texture.texture_descriptor.usage =
-        TextureUsages::COPY_DST | TextureUsages::STORAGE_BINDING | TextureUsages::TEXTURE_BINDING;
-    images.add(output_texture)
-}
-
-impl Plugin for ShocoVoxRenderPlugin {
+impl<T, const DIM: usize> Plugin for RenderBevyPlugin<T, DIM>
+where
+    T: Default + Clone + Copy + PartialEq + VoxelData + Send + Sync + 'static,
+{
     fn build(&self, app: &mut App) {
-        app.add_plugins(ExtractResourcePlugin::<ShocoVoxViewingGlass>::default());
-        app.add_plugins(ExtractResourcePlugin::<ShocoVoxRenderData>::default());
+        app.add_plugins((
+            ExtractResourcePlugin::<OctreeGPUHost<T, DIM>>::default(),
+            ExtractResourcePlugin::<SvxViewSet>::default(),
+        ));
         let render_app = app.sub_app_mut(RenderApp);
+        render_app.add_systems(ExtractSchedule, sync_with_main_world);
         render_app.add_systems(
             Render,
-            prepare_bind_groups.in_set(RenderSet::PrepareBindGroups),
+            (
+                write_to_gpu::<T, DIM>.in_set(RenderSet::PrepareResources),
+                prepare_bind_groups::<T, DIM>.in_set(RenderSet::PrepareBindGroups),
+                handle_gpu_readback::<T, DIM>.in_set(RenderSet::Cleanup),
+            ),
         );
-
         let mut render_graph = render_app.world_mut().resource_mut::<RenderGraph>();
         render_graph.add_node(
-            ShocoVoxLabel,
-            ShocoVoxRenderNode {
+            SvxLabel,
+            SvxRenderNode {
                 ready: false,
                 resolution: self.resolution,
             },
         );
-        render_graph.add_node_edge(ShocoVoxLabel, bevy::render::graph::CameraDriverLabel);
+        render_graph.add_node_edge(SvxLabel, bevy::render::graph::CameraDriverLabel);
     }
 
     fn finish(&self, app: &mut App) {
         let render_app = app.sub_app_mut(RenderApp);
-        render_app.init_resource::<ShocoVoxRenderPipeline>();
-    }
-}
-
-const WORKGROUP_SIZE: u32 = 8;
-impl render_graph::Node for ShocoVoxRenderNode {
-    fn update(&mut self, world: &mut World) {
-        let pipeline = world.resource::<ShocoVoxRenderPipeline>();
-        let render_data = world.get_resource::<ShocoVoxRenderData>();
-        let pipeline_cache = world.resource::<PipelineCache>();
-        if !self.ready {
-            if let CachedPipelineState::Ok(_) =
-                pipeline_cache.get_compute_pipeline_state(pipeline.update_pipeline)
-            {
-                self.ready = render_data.is_some();
-            }
-        }
-    }
-
-    fn run(
-        &self,
-        _graph: &mut render_graph::RenderGraphContext,
-        render_context: &mut RenderContext,
-        world: &World,
-    ) -> Result<(), render_graph::NodeRunError> {
-        let pipeline_cache = world.resource::<PipelineCache>();
-        let pipeline = world.resource::<ShocoVoxRenderPipeline>();
-
-        if self.ready {
-            let mut pass = render_context
-                .command_encoder()
-                .begin_compute_pass(&ComputePassDescriptor::default());
-
-            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();
-            pass.set_pipeline(pipeline);
-            pass.dispatch_workgroups(
-                self.resolution[0] / WORKGROUP_SIZE,
-                self.resolution[1] / WORKGROUP_SIZE,
-                1,
-            );
-        }
-        Ok(())
+        render_app.init_resource::<SvxRenderPipeline>();
     }
 }
diff --git a/src/octree/raytracing/bevy/pipeline.rs b/src/octree/raytracing/bevy/pipeline.rs
new file mode 100644
index 0000000..0c1fe4d
--- /dev/null
+++ b/src/octree/raytracing/bevy/pipeline.rs
@@ -0,0 +1,426 @@
+use crate::octree::{
+    raytracing::bevy::types::{
+        OctreeRenderData, OctreeSpyGlass, SvxRenderNode, SvxRenderPipeline, Viewport,
+    },
+    VoxelData,
+};
+
+use bevy::{
+    asset::AssetServer,
+    ecs::{
+        system::{Res, ResMut},
+        world::{FromWorld, World},
+    },
+    render::{
+        render_asset::RenderAssets,
+        render_graph::{self},
+        render_resource::{
+            encase::{StorageBuffer, UniformBuffer},
+            AsBindGroup, BindGroupEntry, BindingResource, BufferDescriptor, BufferInitDescriptor,
+            BufferUsages, CachedPipelineState, ComputePassDescriptor, ComputePipelineDescriptor,
+            PipelineCache, ShaderSize,
+        },
+        renderer::{RenderContext, RenderDevice, RenderQueue},
+        texture::GpuImage,
+    },
+};
+use std::borrow::Cow;
+
+use super::types::{OctreeRenderDataResources, SvxViewSet};
+
+impl FromWorld for SvxRenderPipeline {
+    fn from_world(world: &mut World) -> Self {
+        let render_device = world.resource::<RenderDevice>();
+        let spyglass_bind_group_layout = OctreeSpyGlass::bind_group_layout(render_device);
+        let render_data_bind_group_layout = OctreeRenderData::bind_group_layout(render_device);
+        let shader = world
+            .resource::<AssetServer>()
+            .load("shaders/viewport_render.wgsl");
+        let pipeline_cache = world.resource::<PipelineCache>();
+        let update_pipeline = pipeline_cache.queue_compute_pipeline(ComputePipelineDescriptor {
+            label: None,
+            layout: vec![
+                spyglass_bind_group_layout.clone(),
+                render_data_bind_group_layout.clone(),
+            ],
+            push_constant_ranges: Vec::new(),
+            shader,
+            shader_defs: vec![],
+            entry_point: Cow::from("update"),
+        });
+
+        SvxRenderPipeline {
+            render_queue: world.resource::<RenderQueue>().clone(),
+            update_tree: true,
+            spyglass_bind_group_layout,
+            render_data_bind_group_layout,
+            update_pipeline,
+            resources: None,
+        }
+    }
+}
+
+//##############################################################################
+//  ███████████   █████  █████ ██████   █████
+// ░░███░░░░░███ ░░███  ░░███ ░░██████ ░░███
+//  ░███    ░███  ░███   ░███  ░███░███ ░███
+//  ░██████████   ░███   ░███  ░███░░███░███
+//  ░███░░░░░███  ░███   ░███  ░███ ░░██████
+//  ░███    ░███  ░███   ░███  ░███  ░░█████
+//  █████   █████ ░░████████   █████  ░░█████
+// ░░░░░   ░░░░░   ░░░░░░░░   ░░░░░    ░░░░░
+//##############################################################################
+const WORKGROUP_SIZE: u32 = 8;
+impl render_graph::Node for SvxRenderNode {
+    fn update(&mut self, world: &mut World) {
+        {
+            let svx_pipeline = world.resource::<SvxRenderPipeline>();
+            let pipeline_cache = world.resource::<PipelineCache>();
+            if !self.ready {
+                if let CachedPipelineState::Ok(_) =
+                    pipeline_cache.get_compute_pipeline_state(svx_pipeline.update_pipeline)
+                {
+                    self.ready = !world.resource::<SvxViewSet>().views.is_empty();
+                }
+            }
+        }
+    }
+
+    fn run(
+        &self,
+        _graph: &mut render_graph::RenderGraphContext,
+        render_context: &mut RenderContext,
+        world: &World,
+    ) -> Result<(), render_graph::NodeRunError> {
+        if self.ready {
+            let pipeline_cache = world.resource::<PipelineCache>();
+            let svx_pipeline = world.resource::<SvxRenderPipeline>();
+            let svx_viewset = world.resource::<SvxViewSet>();
+            let current_view = svx_viewset.views[0].lock().unwrap();
+            let command_encoder = render_context.command_encoder();
+            let data_handler = &current_view.data_handler;
+            let resources = svx_pipeline.resources.as_ref().unwrap();
+            {
+                let mut pass =
+                    command_encoder.begin_compute_pass(&ComputePassDescriptor::default());
+
+                pass.set_bind_group(0, &resources.spyglass_bind_group, &[]);
+                pass.set_bind_group(1, &resources.tree_bind_group, &[]);
+                let pipeline = pipeline_cache
+                    .get_compute_pipeline(svx_pipeline.update_pipeline)
+                    .unwrap();
+                pass.set_pipeline(pipeline);
+                pass.dispatch_workgroups(
+                    self.resolution[0] / WORKGROUP_SIZE,
+                    self.resolution[1] / WORKGROUP_SIZE,
+                    1,
+                );
+            }
+
+            command_encoder.copy_buffer_to_buffer(
+                &resources.metadata_buffer,
+                0,
+                &resources.readable_metadata_buffer,
+                0,
+                (std::mem::size_of_val(&data_handler.render_data.metadata[0])
+                    * data_handler.render_data.metadata.len()) as u64,
+            );
+
+            debug_assert!(
+                !current_view.spyglass.node_requests.is_empty(),
+                "Expected node requests array to not be empty"
+            );
+            command_encoder.copy_buffer_to_buffer(
+                &resources.node_requests_buffer,
+                0,
+                &resources.readable_node_requests_buffer,
+                0,
+                (std::mem::size_of_val(&current_view.spyglass.node_requests[0])
+                    * current_view.spyglass.node_requests.len()) as u64,
+            );
+        }
+        Ok(())
+    }
+}
+
+//##############################################################################
+//    █████████  ███████████   ██████████   █████████   ███████████ ██████████
+//   ███░░░░░███░░███░░░░░███ ░░███░░░░░█  ███░░░░░███ ░█░░░███░░░█░░███░░░░░█
+//  ███     ░░░  ░███    ░███  ░███  █ ░  ░███    ░███ ░   ░███  ░  ░███  █ ░
+// ░███          ░██████████   ░██████    ░███████████     ░███     ░██████
+// ░███          ░███░░░░░███  ░███░░█    ░███░░░░░███     ░███     ░███░░█
+// ░░███     ███ ░███    ░███  ░███ ░   █ ░███    ░███     ░███     ░███ ░   █
+//  ░░█████████  █████   █████ ██████████ █████   █████    █████    ██████████
+//   ░░░░░░░░░  ░░░░░   ░░░░░ ░░░░░░░░░░ ░░░░░   ░░░░░    ░░░░░    ░░░░░░░░░░
+//  ███████████  █████ ██████   █████ ██████████
+// ░░███░░░░░███░░███ ░░██████ ░░███ ░░███░░░░███
+//  ░███    ░███ ░███  ░███░███ ░███  ░███   ░░███
+//  ░██████████  ░███  ░███░░███░███  ░███    ░███
+//  ░███░░░░░███ ░███  ░███ ░░██████  ░███    ░███
+//  ░███    ░███ ░███  ░███  ░░█████  ░███    ███
+//  ███████████  █████ █████  ░░█████ ██████████
+// ░░░░░░░░░░░  ░░░░░ ░░░░░    ░░░░░ ░░░░░░░░░░
+//    █████████  ███████████      ███████    █████  █████ ███████████   █████████
+//   ███░░░░░███░░███░░░░░███   ███░░░░░███ ░░███  ░░███ ░░███░░░░░███ ███░░░░░███
+//  ███     ░░░  ░███    ░███  ███     ░░███ ░███   ░███  ░███    ░███░███    ░░░
+// ░███          ░██████████  ░███      ░███ ░███   ░███  ░██████████ ░░█████████
+// ░███    █████ ░███░░░░░███ ░███      ░███ ░███   ░███  ░███░░░░░░   ░░░░░░░░███
+// ░░███  ░░███  ░███    ░███ ░░███     ███  ░███   ░███  ░███         ███    ░███
+//  ░░█████████  █████   █████ ░░░███████░   ░░████████   █████       ░░█████████
+//   ░░░░░░░░░  ░░░░░   ░░░░░    ░░░░░░░      ░░░░░░░░   ░░░░░         ░░░░░░░░░
+//##############################################################################
+/// Constructs buffers, bing groups and uploads rendering data at initialization and whenever prompted
+pub(crate) fn prepare_bind_groups<T, const DIM: usize>(
+    gpu_images: Res<RenderAssets<GpuImage>>,
+    render_device: Res<RenderDevice>,
+    mut pipeline: ResMut<SvxRenderPipeline>,
+    svx_viewset: ResMut<SvxViewSet>,
+) where
+    T: Default + Clone + PartialEq + VoxelData + Send + Sync + 'static,
+{
+    if pipeline.resources.is_some() && !pipeline.update_tree {
+        return;
+    }
+
+    let tree_view = &svx_viewset.views[0].lock().unwrap();
+    let render_data = &tree_view.data_handler.render_data;
+    if let Some(resources) = &pipeline.resources {
+        let mut buffer = UniformBuffer::new(Vec::<u8>::new());
+        buffer.write(&render_data.octree_meta).unwrap();
+        pipeline
+            .render_queue
+            .write_buffer(&resources.metadata_buffer, 0, &buffer.into_inner());
+
+        let mut buffer = StorageBuffer::new(Vec::<u8>::new());
+        buffer.write(&render_data.metadata).unwrap();
+        pipeline
+            .render_queue
+            .write_buffer(&resources.metadata_buffer, 0, &buffer.into_inner());
+
+        let mut buffer = StorageBuffer::new(Vec::<u8>::new());
+        buffer.write(&render_data.node_children).unwrap();
+        pipeline.render_queue.write_buffer(
+            &resources.node_children_buffer,
+            0,
+            &buffer.into_inner(),
+        );
+
+        let mut buffer = StorageBuffer::new(Vec::<u8>::new());
+        buffer.write(&render_data.node_ocbits).unwrap();
+        pipeline
+            .render_queue
+            .write_buffer(&resources.node_ocbits_buffer, 0, &buffer.into_inner());
+
+        let mut buffer = StorageBuffer::new(Vec::<u8>::new());
+        buffer.write(&render_data.voxels).unwrap();
+        pipeline
+            .render_queue
+            .write_buffer(&resources.voxels_buffer, 0, &buffer.into_inner());
+
+        let mut buffer = StorageBuffer::new(Vec::<u8>::new());
+        buffer.write(&render_data.color_palette).unwrap();
+        pipeline
+            .render_queue
+            .write_buffer(&resources.color_palette_buffer, 0, &buffer.into_inner())
+    } else {
+        //##############################################################################
+        //  ███████████ ███████████   ██████████ ██████████
+        // ░█░░░███░░░█░░███░░░░░███ ░░███░░░░░█░░███░░░░░█
+        // ░   ░███  ░  ░███    ░███  ░███  █ ░  ░███  █ ░
+        //     ░███     ░██████████   ░██████    ░██████
+        //     ░███     ░███░░░░░███  ░███░░█    ░███░░█
+        //     ░███     ░███    ░███  ░███ ░   █ ░███ ░   █
+        //     █████    █████   █████ ██████████ ██████████
+        //    ░░░░░    ░░░░░   ░░░░░ ░░░░░░░░░░ ░░░░░░░░░░
+        //    █████████  ███████████      ███████    █████  █████ ███████████
+        //   ███░░░░░███░░███░░░░░███   ███░░░░░███ ░░███  ░░███ ░░███░░░░░███
+        //  ███     ░░░  ░███    ░███  ███     ░░███ ░███   ░███  ░███    ░███
+        // ░███          ░██████████  ░███      ░███ ░███   ░███  ░██████████
+        // ░███    █████ ░███░░░░░███ ░███      ░███ ░███   ░███  ░███░░░░░░
+        // ░░███  ░░███  ░███    ░███ ░░███     ███  ░███   ░███  ░███
+        //  ░░█████████  █████   █████ ░░░███████░   ░░████████   █████
+        //   ░░░░░░░░░  ░░░░░   ░░░░░    ░░░░░░░      ░░░░░░░░   ░░░░░
+        //##############################################################################
+        // Create the staging buffer helping in reading data from the GPU
+        let readable_metadata_buffer = render_device.create_buffer(&BufferDescriptor {
+            mapped_at_creation: false,
+            size: (render_data.metadata.len() * 4) as u64,
+            label: Some("Octree Node metadata staging Buffer"),
+            usage: BufferUsages::COPY_DST | BufferUsages::MAP_READ,
+        });
+
+        let mut buffer = UniformBuffer::new(Vec::<u8>::new());
+        buffer.write(&render_data.octree_meta).unwrap();
+        let octree_meta_buffer = render_device.create_buffer_with_data(&BufferInitDescriptor {
+            label: Some("Octree Tree Metadata Buffer"),
+            contents: &buffer.into_inner(),
+            usage: BufferUsages::UNIFORM | BufferUsages::COPY_DST,
+        });
+
+        let mut buffer = StorageBuffer::new(Vec::<u8>::new());
+        buffer.write(&render_data.metadata).unwrap();
+        let metadata_buffer = render_device.create_buffer_with_data(&BufferInitDescriptor {
+            label: Some("Octree Metadata Buffer"),
+            contents: &buffer.into_inner(),
+            usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC | BufferUsages::COPY_DST,
+        });
+
+        let mut buffer = StorageBuffer::new(Vec::<u8>::new());
+        buffer.write(&render_data.node_children).unwrap();
+        let node_children_buffer = render_device.create_buffer_with_data(&BufferInitDescriptor {
+            label: Some("Octree Node Children Buffer"),
+            contents: &buffer.into_inner(),
+            usage: BufferUsages::STORAGE | BufferUsages::COPY_DST,
+        });
+
+        let mut buffer = StorageBuffer::new(Vec::<u8>::new());
+        buffer.write(&render_data.node_ocbits).unwrap();
+        let node_ocbits_buffer = render_device.create_buffer_with_data(&BufferInitDescriptor {
+            label: Some("Octree Node Occupied Bits Buffer"),
+            contents: &buffer.into_inner(),
+            usage: BufferUsages::STORAGE | BufferUsages::COPY_DST,
+        });
+
+        let mut buffer = StorageBuffer::new(Vec::<u8>::new());
+        buffer.write(&render_data.voxels).unwrap();
+        let voxels_buffer = render_device.create_buffer_with_data(&BufferInitDescriptor {
+            label: Some("Octree Voxels Buffer"),
+            contents: &buffer.into_inner(),
+            usage: BufferUsages::STORAGE | BufferUsages::COPY_DST,
+        });
+
+        let mut buffer = StorageBuffer::new(Vec::<u8>::new());
+        buffer.write(&render_data.color_palette).unwrap();
+        let color_palette_buffer = render_device.create_buffer_with_data(&BufferInitDescriptor {
+            label: Some("Octree Color Palette Buffer"),
+            contents: &buffer.into_inner(),
+            usage: BufferUsages::STORAGE | BufferUsages::COPY_DST,
+        });
+
+        // Create bind group
+        let tree_bind_group = render_device.create_bind_group(
+            OctreeRenderData::label(),
+            &pipeline.render_data_bind_group_layout,
+            &[
+                bevy::render::render_resource::BindGroupEntry {
+                    binding: 0,
+                    resource: octree_meta_buffer.as_entire_binding(),
+                },
+                bevy::render::render_resource::BindGroupEntry {
+                    binding: 1,
+                    resource: metadata_buffer.as_entire_binding(),
+                },
+                bevy::render::render_resource::BindGroupEntry {
+                    binding: 2,
+                    resource: node_children_buffer.as_entire_binding(),
+                },
+                bevy::render::render_resource::BindGroupEntry {
+                    binding: 3,
+                    resource: node_ocbits_buffer.as_entire_binding(),
+                },
+                bevy::render::render_resource::BindGroupEntry {
+                    binding: 4,
+                    resource: voxels_buffer.as_entire_binding(),
+                },
+                bevy::render::render_resource::BindGroupEntry {
+                    binding: 5,
+                    resource: color_palette_buffer.as_entire_binding(),
+                },
+            ],
+        );
+
+        //##############################################################################
+        //   █████████  ███████████  █████ █████
+        //  ███░░░░░███░░███░░░░░███░░███ ░░███
+        // ░███    ░░░  ░███    ░███ ░░███ ███
+        // ░░█████████  ░██████████   ░░█████
+        //  ░░░░░░░░███ ░███░░░░░░     ░░███
+        //  ███    ░███ ░███            ░███
+        // ░░█████████  █████           █████
+        //  ░░░░░░░░░  ░░░░░           ░░░░░
+        //    █████████  █████         █████████    █████████   █████████
+        //   ███░░░░░███░░███         ███░░░░░███  ███░░░░░███ ███░░░░░███
+        //  ███     ░░░  ░███        ░███    ░███ ░███    ░░░ ░███    ░░░
+        // ░███          ░███        ░███████████ ░░█████████ ░░█████████
+        // ░███    █████ ░███        ░███░░░░░███  ░░░░░░░░███ ░░░░░░░░███
+        // ░░███  ░░███  ░███      █ ░███    ░███  ███    ░███ ███    ░███
+        //  ░░█████████  ███████████ █████   █████░░█████████ ░░█████████
+        //   ░░░░░░░░░  ░░░░░░░░░░░ ░░░░░   ░░░░░  ░░░░░░░░░   ░░░░░░░░░
+        //    █████████  ███████████      ███████    █████  █████ ███████████
+        //   ███░░░░░███░░███░░░░░███   ███░░░░░███ ░░███  ░░███ ░░███░░░░░███
+        //  ███     ░░░  ░███    ░███  ███     ░░███ ░███   ░███  ░███    ░███
+        // ░███          ░██████████  ░███      ░███ ░███   ░███  ░██████████
+        // ░███    █████ ░███░░░░░███ ░███      ░███ ░███   ░███  ░███░░░░░░
+        // ░░███  ░░███  ░███    ░███ ░░███     ███  ░███   ░███  ░███
+        //  ░░█████████  █████   █████ ░░░███████░   ░░████████   █████
+        //##############################################################################
+        let mut buffer = UniformBuffer::new([0u8; Viewport::SHADER_SIZE.get() as usize]);
+        buffer.write(&tree_view.spyglass.viewport).unwrap();
+        let viewport_buffer = render_device.create_buffer_with_data(&BufferInitDescriptor {
+            label: Some("Octree Viewport Buffer"),
+            contents: &buffer.into_inner(),
+            usage: BufferUsages::UNIFORM | BufferUsages::COPY_DST,
+        });
+
+        debug_assert!(
+            !tree_view.spyglass.node_requests.is_empty(),
+            "Expected node requests array to not be empty"
+        );
+        let mut buffer = StorageBuffer::new(Vec::<u8>::new());
+        buffer.write(&tree_view.spyglass.node_requests).unwrap();
+        let node_requests_buffer = render_device.create_buffer_with_data(&BufferInitDescriptor {
+            label: Some("Octree Node requests Buffer"),
+            contents: &buffer.into_inner(),
+            usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC | BufferUsages::COPY_DST,
+        });
+
+        let readable_node_requests_buffer = render_device.create_buffer(&BufferDescriptor {
+            mapped_at_creation: false,
+            size: (tree_view.spyglass.node_requests.len()
+                * std::mem::size_of_val(&tree_view.spyglass.node_requests[0]))
+                as u64,
+            label: Some("Octree Node requests staging Buffer"),
+            usage: BufferUsages::COPY_DST | BufferUsages::MAP_READ,
+        });
+
+        let output_texture_view = gpu_images
+            .get(&tree_view.spyglass.output_texture)
+            .unwrap()
+            .texture_view
+            .clone();
+        let spyglass_bind_group = render_device.create_bind_group(
+            OctreeSpyGlass::label(),
+            &pipeline.spyglass_bind_group_layout,
+            &[
+                BindGroupEntry {
+                    binding: 0,
+                    resource: BindingResource::TextureView(&output_texture_view.clone()),
+                },
+                BindGroupEntry {
+                    binding: 1,
+                    resource: viewport_buffer.as_entire_binding(),
+                },
+                BindGroupEntry {
+                    binding: 2,
+                    resource: node_requests_buffer.as_entire_binding(),
+                },
+            ],
+        );
+
+        pipeline.resources = Some(OctreeRenderDataResources {
+            node_requests_buffer,
+            spyglass_bind_group,
+            tree_bind_group,
+            viewport_buffer,
+            metadata_buffer,
+            node_children_buffer,
+            node_ocbits_buffer,
+            voxels_buffer,
+            color_palette_buffer,
+            readable_node_requests_buffer,
+            readable_metadata_buffer,
+        });
+    }
+
+    pipeline.update_tree = false;
+}
diff --git a/src/octree/raytracing/bevy/types.rs b/src/octree/raytracing/bevy/types.rs
index 94049fa..5841190 100644
--- a/src/octree/raytracing/bevy/types.rs
+++ b/src/octree/raytracing/bevy/types.rs
@@ -1,4 +1,4 @@
-use crate::octree::V3cf32;
+use crate::octree::{Albedo, Octree, V3cf32, VoxelData};
 use bevy::{
     asset::Handle,
     ecs::system::Resource,
@@ -9,10 +9,16 @@ use bevy::{
         prelude::Image,
         render_graph::RenderLabel,
         render_resource::{
-            AsBindGroup, BindGroup, BindGroupLayout, CachedComputePipelineId, ShaderType,
+            AsBindGroup, BindGroup, BindGroupLayout, Buffer, CachedComputePipelineId, ShaderType,
         },
+        renderer::RenderQueue,
     },
 };
+use bimap::BiHashMap;
+use std::{
+    collections::HashMap,
+    sync::{Arc, Mutex},
+};
 
 #[derive(Clone, ShaderType)]
 pub(crate) struct Voxelement {
@@ -35,33 +41,107 @@ pub struct Viewport {
     pub w_h_fov: V3cf32,
 }
 
-pub struct ShocoVoxRenderPlugin {
-    pub resolution: [u32; 2],
+pub struct RenderBevyPlugin<T, const DIM: usize>
+where
+    T: Default + Clone + PartialEq + VoxelData + Send + Sync + 'static,
+{
+    pub(crate) dummy: std::marker::PhantomData<T>,
+    pub(crate) resolution: [u32; 2],
 }
 
-#[derive(Resource, Clone, AsBindGroup, TypePath, ExtractResource)]
-#[type_path = "shocovox::gpu::ShocoVoxViewingGlass"]
-pub struct ShocoVoxViewingGlass {
+#[derive(Resource, Clone, TypePath, ExtractResource)]
+#[type_path = "shocovox::gpu::OctreeGPUHost"]
+pub struct OctreeGPUHost<T, const DIM: usize>
+where
+    T: Default + Clone + PartialEq + VoxelData + Send + Sync + 'static,
+{
+    pub tree: Octree<T, DIM>,
+}
+
+#[derive(Default, Resource, Clone, TypePath, ExtractResource)]
+#[type_path = "shocovox::gpu::SvxViewSet"]
+pub struct SvxViewSet {
+    pub views: Vec<Arc<Mutex<OctreeGPUView>>>,
+}
+
+#[derive(Resource, Clone, AsBindGroup)]
+pub struct OctreeGPUView {
+    pub spyglass: OctreeSpyGlass,
+    pub(crate) data_handler: OctreeGPUDataHandler,
+}
+
+#[derive(Debug, Clone)]
+pub(crate) struct VictimPointer {
+    pub(crate) max_meta_len: usize,
+    pub(crate) loop_count: usize,
+    pub(crate) stored_items: usize,
+    pub(crate) meta_index: usize,
+    pub(crate) child: usize,
+}
+
+#[derive(Debug, Clone, PartialEq)]
+pub(crate) enum BrickOwnedBy {
+    NotOwned,
+    Node(u32, u8),
+}
+
+#[derive(Resource, Clone, AsBindGroup)]
+pub struct OctreeGPUDataHandler {
+    pub(crate) render_data: OctreeRenderData,
+    pub(crate) victim_node: VictimPointer,
+    pub(crate) victim_brick: usize,
+    pub(crate) node_key_vs_meta_index: BiHashMap<usize, usize>,
+    pub(crate) map_to_color_index_in_palette: HashMap<Albedo, usize>,
+    pub(crate) brick_ownership: Vec<BrickOwnedBy>,
+    pub(crate) map_to_brick_maybe_owned_by_node: HashMap<(usize, u8), usize>,
+    pub(crate) uploaded_color_palette_size: usize,
+}
+
+#[derive(Clone)]
+pub(crate) struct OctreeRenderDataResources {
+    // Spyglass group
+    pub(crate) spyglass_bind_group: BindGroup,
+    pub(crate) viewport_buffer: Buffer,
+    pub(crate) node_requests_buffer: Buffer,
+
+    // Octree render data group
+    pub(crate) tree_bind_group: BindGroup,
+    pub(crate) metadata_buffer: Buffer,
+    pub(crate) node_children_buffer: Buffer,
+    pub(crate) node_ocbits_buffer: Buffer,
+    pub(crate) voxels_buffer: Buffer,
+    pub(crate) color_palette_buffer: Buffer,
+
+    // Staging buffers for data reads
+    pub(crate) readable_node_requests_buffer: Buffer,
+    pub(crate) readable_metadata_buffer: Buffer,
+}
+
+#[derive(Clone, AsBindGroup)]
+pub struct OctreeSpyGlass {
     #[storage_texture(0, image_format = Rgba8Unorm, access = ReadWrite)]
     pub output_texture: Handle<Image>,
 
     #[uniform(1, visibility(compute))]
     pub viewport: Viewport,
+
+    #[storage(2, visibility(compute))]
+    pub(crate) node_requests: Vec<u32>,
 }
 
-#[derive(Resource, Clone, AsBindGroup, TypePath, ExtractResource)]
+#[derive(Clone, AsBindGroup, TypePath)]
 #[type_path = "shocovox::gpu::ShocoVoxRenderData"]
-pub struct ShocoVoxRenderData {
+pub struct OctreeRenderData {
+    /// Contains the properties of the Octree
     #[uniform(0, visibility(compute))]
-    pub(crate) meta: OctreeMetaData,
+    pub(crate) octree_meta: OctreeMetaData,
 
-    /// Composite field containing the properties of Nodes
-    /// Structure is the following:
+    /// Contains the properties of Nodes and Voxel Bricks
     ///  _===================================================================_
     /// | Byte 0  | Node properties                                           |
     /// |---------------------------------------------------------------------|
-    /// |  bit 0  | unused - potentially: "node in use do not delete" bit     |
-    /// |  bit 1  | unused - potentially: "brick in use do not delete" bit    |
+    /// |  bit 0  | 1 if node is used by the raytracing algorithm *(2) *(4)   |
+    /// |  bit 1  | unused                                                    |
     /// |  bit 2  | 1 in case node is a leaf                                  |
     /// |  bit 3  | 1 in case node is uniform                                 |
     /// |  bit 4  | unused - potentially: 1 if node has voxels                |
@@ -69,21 +149,30 @@ pub struct ShocoVoxRenderData {
     /// |  bit 6  | unused - potentially: voxel brick size: 1, full or sparse |
     /// |  bit 7  | unused                                                    |
     /// |=====================================================================|
-    /// | Byte 1  | Child occupied                                            |
+    /// | Byte 1  | Node Child occupied                                       |
     /// |---------------------------------------------------------------------|
     /// | If Leaf | each bit is 0 if child brick is empty at octant *(1)      |
     /// | If Node | unused                                                    |
     /// |=====================================================================|
-    /// | Byte 2  | Child structure                                           |
+    /// | Byte 2  | Node Child structure                                      |
     /// |---------------------------------------------------------------------|
     /// | If Leaf | each bit is 0 if child brick is solid, 1 if parted *(1)   |
     /// | If Node | unused                                                    |
     /// |=====================================================================|
-    /// | Byte 3  | unused                                                    |
+    /// | Byte 3  | Voxel Bricks used *(3)                                    |
+    /// |---------------------------------------------------------------------|
+    /// | each bit is 1 if brick is used (do not delete please)               |
     /// `=====================================================================`
-    /// *(1) Only first bit is used in case leaf is uniform
+    /// *(1) Only first bit is used in case uniform leaf nodes
+    /// *(2) The same bit is used for node_children and node_occupied_bits
+    /// *(3) One index in the array covers 8 bricks, which is the theoretical maximum
+    ///      number of bricks for one node. In practice however the number of bricks
+    ///      are only 4-5 times more, than the number of nodes, because of the internal nodes;
+    ///      And only a fraction of them are visible in a render.
+    /// *(4) Root node does not have this bit used, because it will never be overwritten
+    ///      due to the victim pointer logic
     #[storage(1, visibility(compute))]
-    pub(crate) nodes: Vec<u32>,
+    pub(crate) metadata: Vec<u32>,
 
     /// Index values for Nodes, 8 value per @SizedNode entry. Each value points to:
     /// In case of Internal Nodes
@@ -94,17 +183,17 @@ pub struct ShocoVoxRenderData {
     /// index of where the voxel brick start inside the @voxels buffer.
     /// Leaf node might contain 1 or 8 bricks according to @sized_node_meta, while
     #[storage(2, visibility(compute))]
-    pub(crate) children_buffer: Vec<u32>,
-
-    /// Buffer of Voxel Bricks. Each brick contains voxel_brick_dim^3 elements.
-    /// Each Brick has a corresponding 64 bit occupancy bitmap in the @voxel_maps buffer.
-    #[storage(3, visibility(compute))]
-    pub(crate) voxels: Vec<Voxelement>,
+    pub(crate) node_children: Vec<u32>,
 
     /// Buffer of Node occupancy bitmaps. Each node has a 64 bit bitmap,
     /// which is stored in 2 * u32 values
+    #[storage(3, visibility(compute))]
+    pub(crate) node_ocbits: Vec<u32>,
+
+    /// Buffer of Voxel Bricks. Each brick contains voxel_brick_dim^3 elements.
+    /// Each Brick has a corresponding 64 bit occupancy bitmap in the @voxel_maps buffer.
     #[storage(4, visibility(compute))]
-    pub(crate) node_occupied_bits: Vec<u32>,
+    pub(crate) voxels: Vec<Voxelement>,
 
     /// Stores each unique color, it is references in @voxels
     /// and in @children_buffer as well( in case of solid bricks )
@@ -113,19 +202,22 @@ pub struct ShocoVoxRenderData {
 }
 
 #[derive(Resource)]
-pub(crate) struct ShocoVoxRenderPipeline {
+pub(crate) struct SvxRenderPipeline {
     pub update_tree: bool,
-    pub(crate) viewing_glass_bind_group_layout: BindGroupLayout,
-    pub(crate) render_data_bind_group_layout: BindGroupLayout,
+
+    pub(crate) render_queue: RenderQueue,
     pub(crate) update_pipeline: CachedComputePipelineId,
-    pub(crate) viewing_glass_bind_group: Option<BindGroup>,
-    pub(crate) tree_bind_group: Option<BindGroup>,
+
+    // Data layout and data
+    pub(crate) spyglass_bind_group_layout: BindGroupLayout,
+    pub(crate) render_data_bind_group_layout: BindGroupLayout,
+    pub(crate) resources: Option<OctreeRenderDataResources>,
 }
 
 #[derive(Debug, Hash, PartialEq, Eq, Clone, RenderLabel)]
-pub(crate) struct ShocoVoxLabel;
+pub(crate) struct SvxLabel;
 
-pub(crate) struct ShocoVoxRenderNode {
+pub(crate) struct SvxRenderNode {
     pub(crate) ready: bool,
     pub(crate) resolution: [u32; 2],
 }
diff --git a/src/octree/raytracing/mod.rs b/src/octree/raytracing/mod.rs
index f54d278..66e9ec7 100644
--- a/src/octree/raytracing/mod.rs
+++ b/src/octree/raytracing/mod.rs
@@ -7,4 +7,7 @@ pub mod bevy;
 pub use crate::spatial::raytracing::Ray;
 
 #[cfg(feature = "bevy_wgpu")]
-pub use bevy::types::{ShocoVoxRenderPlugin, ShocoVoxViewingGlass, Viewport};
+pub use bevy::types::{
+    OctreeGPUHost, OctreeGPUView, OctreeRenderData, OctreeSpyGlass, RenderBevyPlugin, SvxViewSet,
+    Viewport,
+};
diff --git a/src/octree/raytracing/raytracing_on_cpu.rs b/src/octree/raytracing/raytracing_on_cpu.rs
index 03ab521..b1e5e3a 100644
--- a/src/octree/raytracing/raytracing_on_cpu.rs
+++ b/src/octree/raytracing/raytracing_on_cpu.rs
@@ -8,7 +8,7 @@ use crate::{
             BITMAP_INDEX_LUT, BITMAP_MASK_FOR_OCTANT_LUT, OOB_OCTANT,
             RAY_TO_NODE_OCCUPANCY_BITMASK_LUT,
         },
-        math::{hash_direction, hash_region, position_in_bitmap_64bits, BITMAP_DIMENSION},
+        math::{hash_direction, hash_region, BITMAP_DIMENSION},
         raytracing::{cube_impact_normal, step_octant, Ray, FLOAT_ERROR_TOLERANCE},
     },
 };
@@ -294,9 +294,9 @@ where
         let mut step_vec = V3c::unit(0.);
 
         while target_octant != OOB_OCTANT {
-            current_node_key = Octree::<T, DIM>::ROOT_NODE_KEY as usize;
+            current_node_key = Self::ROOT_NODE_KEY as usize;
             current_bounds = Cube::root_bounds(self.octree_size as f32);
-            node_stack.push(Octree::<T, DIM>::ROOT_NODE_KEY);
+            node_stack.push(Self::ROOT_NODE_KEY);
             while !node_stack.is_empty() {
                 let current_node_occupied_bits =
                     self.stored_occupied_bits(*node_stack.last().unwrap() as usize);
diff --git a/src/octree/types.rs b/src/octree/types.rs
index 2b05283..ebebfdd 100644
--- a/src/octree/types.rs
+++ b/src/octree/types.rs
@@ -77,6 +77,7 @@ pub trait VoxelData {
 /// A Brick can be indexed directly, as opposed to the octree which is essentially a
 /// tree-graph where each node has 8 children.
 #[cfg_attr(feature = "serialization", derive(Serialize))]
+#[derive(Clone)]
 pub struct Octree<T, const DIM: usize = 1>
 where
     T: Default + Clone + PartialEq + VoxelData,
diff --git a/src/octree/update.rs b/src/octree/update.rs
index 61b1a4e..0e7ba37 100644
--- a/src/octree/update.rs
+++ b/src/octree/update.rs
@@ -332,7 +332,7 @@ where
         }
 
         // A CPU stack does not consume significant relevant resources, e.g. a 4096*4096*4096 chunk has depth of 12
-        let mut node_stack = vec![(Octree::<T, DIM>::ROOT_NODE_KEY, root_bounds)];
+        let mut node_stack = vec![(Self::ROOT_NODE_KEY, root_bounds)];
         let mut actual_update_size = 0;
         loop {
             let (current_node_key, current_bounds) = *node_stack.last().unwrap();
@@ -472,8 +472,8 @@ where
             if let NodeContent::Internal(ref mut occupied_bits) =
                 self.nodes.get_mut(node_key as usize)
             {
-                let corrected_update_size = ((node_bounds.size * actual_update_size as f32)
-                    / BITMAP_DIMENSION as f32)
+                let corrected_update_size = ((actual_update_size as f32 * BITMAP_DIMENSION as f32)
+                    / node_bounds.size)
                     .ceil() as usize;
                 set_occupancy_in_bitmap_64bits(
                     &matrix_index_for(&node_bounds, &(position.into()), BITMAP_DIMENSION),
@@ -541,7 +541,7 @@ where
         }
 
         // A CPU stack does not consume significant relevant resources, e.g. a 4096*4096*4096 chunk has depth of 12
-        let mut node_stack = vec![(Octree::<T, DIM>::ROOT_NODE_KEY, root_bounds)];
+        let mut node_stack = vec![(Self::ROOT_NODE_KEY, root_bounds)];
         let mut actual_update_size = 0;
         loop {
             let (current_node_key, current_bounds) = *node_stack.last().unwrap();
diff --git a/src/spatial/lut.rs b/src/spatial/lut.rs
index dca996a..da40265 100644
--- a/src/spatial/lut.rs
+++ b/src/spatial/lut.rs
@@ -198,13 +198,13 @@ pub(crate) const OCTANT_OFFSET_REGION_LUT: [V3cf32; 8] = [
 
 pub(crate) const BITMAP_MASK_FOR_OCTANT_LUT: [u64; 8] = [
     0x0000000000330033,
-    0x0000000000cc00cc,
+    0x0000000000CC00CC,
     0x0033003300000000,
-    0x00cc00cc00000000,
+    0x00CC00CC00000000,
     0x0000000033003300,
-    0x00000000cc00cc00,
+    0x00000000CC00CC00,
     0x3300330000000000,
-    0xcc00cc0000000000,
+    0xCC00CC0000000000,
 ];
 
 pub(crate) const BITMAP_INDEX_LUT: [[[usize; 4]; 4]; 4] = [