Add compute worker for visible chunk culling

client/assets/shaders/visible_chunks.wgsl

@@ -0,0 +1,30 @@
+struct Params {
+    centre: vec3<i32>;
+    radius: i32;
+    count: u32;
+    _pad: u32;
+};
+
+@group(0) @binding(0)
+var<uniform> params: Params;
+
+@group(0) @binding(1)
+var<storage, read> keys_in: array<vec3<i32>>;
+
+@group(0) @binding(2)
+var<storage, read_write> results: array<vec4<i32>>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) id: vec3<u32>) {
+    let idx = id.x;
+    if (idx >= params.count) { return; }
+    let key = keys_in[idx];
+    let dx = key.x - params.centre.x;
+    let dy = key.y - params.centre.y;
+    let dz = key.z - params.centre.z;
+    var dist2 = dx * dx + dy * dy + dz * dz;
+    if (abs(dx) > params.radius || abs(dy) > params.radius || abs(dz) > params.radius) {
+        dist2 = -1;
+    }
+    results[idx] = vec4<i32>(key, dist2);
+}

client/src/app.rs

@@ -2,6 +2,7 @@ use bevy::pbr::wireframe::WireframePlugin;
 use crate::helper::debug_gizmos::debug_gizmos;
 use bevy_easy_compute::prelude::{AppComputePlugin, AppComputeWorkerPlugin};
 use crate::plugins::environment::systems::voxels::sphere_compute::SphereWorker;
+use crate::plugins::environment::systems::voxels::visible_chunks_compute::VisibleChunksWorker;
 use bevy::prelude::*;
 pub struct AppPlugin;
 
@@ -12,6 +13,7 @@ impl Plugin for AppPlugin {
         app.add_plugins(crate::plugins::environment::environment_plugin::EnvironmentPlugin);
         app.add_plugins(AppComputePlugin);
         app.add_plugins(AppComputeWorkerPlugin::<SphereWorker>::default());
+        app.add_plugins(AppComputeWorkerPlugin::<VisibleChunksWorker>::default());
         //app.add_plugins(crate::plugins::network::network_plugin::NetworkPlugin);
         app.add_plugins(crate::plugins::input::input_plugin::InputPlugin);
         app.add_plugins(WireframePlugin);

client/src/plugins/environment/environment_plugin.rs

@@ -4,10 +4,11 @@ use bevy_easy_compute::prelude::*;
 use crate::plugins::environment::systems::voxels::sphere_compute::{SphereWorker, SphereParams, SphereGenerated, execute_sphere_once, apply_sphere_result};
 use crate::plugins::environment::systems::voxels::debug::{draw_grid, visualize_octree_system};
 use crate::plugins::environment::systems::voxels::queue_systems;
-use crate::plugins::environment::systems::voxels::queue_systems::{enqueue_visible_chunks, process_chunk_queue};
+use crate::plugins::environment::systems::voxels::queue_systems::{enqueue_visible_chunks, process_chunk_queue, apply_visible_chunk_results};
 use crate::plugins::environment::systems::voxels::render_chunks::rebuild_dirty_chunks;
 use crate::plugins::environment::systems::voxels::lod::update_chunk_lods;
 use crate::plugins::environment::systems::voxels::structure::{ChunkBudget, ChunkCullingCfg, ChunkQueue, SparseVoxelOctree, SpawnedChunks, PrevCameraChunk};
+use crate::plugins::environment::systems::voxels::visible_chunks_compute::VisibleChunkCount;
 
 pub struct EnvironmentPlugin;
 impl Plugin for EnvironmentPlugin {
@@ -35,6 +36,7 @@ impl Plugin for EnvironmentPlugin {
             // ------------------------------------------------------------------------
             .init_resource::<ChunkQueue>()
             .init_resource::<SpawnedChunks>()
+            .init_resource::<VisibleChunkCount>()
             // ------------------------------------------------------------------------
             // frame update
             // ------------------------------------------------------------------------
@@ -43,7 +45,8 @@ impl Plugin for EnvironmentPlugin {
                 (
                     /* ---------- culling & streaming ---------- */
                     enqueue_visible_chunks,
-                    process_chunk_queue.after(enqueue_visible_chunks),
+                    apply_visible_chunk_results.after(enqueue_visible_chunks),
+                    process_chunk_queue.after(apply_visible_chunk_results),
                     update_chunk_lods.after(process_chunk_queue),
                     rebuild_dirty_chunks .after(process_chunk_queue),
                     apply_sphere_result.after(rebuild_dirty_chunks),             // 4.  (re)mesh dirty chunks

client/src/plugins/environment/systems/voxels/mod.rs

@@ -11,3 +11,4 @@ pub mod queue_systems;
 pub mod lod;
 pub mod noise_compute;
 pub mod sphere_compute;
+pub mod visible_chunks_compute;

client/src/plugins/environment/systems/voxels/queue_systems.rs

@@ -1,51 +1,51 @@
 use bevy::prelude::*;
 use crate::plugins::environment::systems::voxels::helper::world_to_chunk;
 use crate::plugins::environment::systems::voxels::structure::*;
+use crate::plugins::environment::systems::voxels::visible_chunks_compute::{
+    VisibleChunksWorker, VisibleParams, IVec3Pod, ChunkResult, VisibleChunkCount,
+    MAX_VISIBLE_CHUNKS,
+};
+use bevy_easy_compute::prelude::*;
 
 
 
 /// enqueue chunks that *should* be visible but are not yet spawned
 /// enqueue chunks that *should* be visible but are not yet spawned
 pub fn enqueue_visible_chunks(
-    mut queue      : ResMut<ChunkQueue>,
-    spawned        : Res<SpawnedChunks>,
     mut prev_cam   : ResMut<PrevCameraChunk>,
     cfg            : Res<ChunkCullingCfg>,
     cam_q          : Query<&GlobalTransform, With<Camera>>,
     tree_q         : Query<&SparseVoxelOctree>,
+    mut worker     : ResMut<AppComputeWorker<VisibleChunksWorker>>,
+    mut count_res  : ResMut<VisibleChunkCount>,
 ) {
-    let tree     = tree_q.single();
-    let cam_pos  = cam_q.single().translation();
-    let centre   = world_to_chunk(tree, cam_pos);
+    let tree    = tree_q.single();
+    let cam_pos = cam_q.single().translation();
+    let centre  = world_to_chunk(tree, cam_pos);
 
     if prev_cam.0 == Some(centre) {
         return;
     }
     prev_cam.0 = Some(centre);
 
-    let r = cfg.view_distance_chunks;
-
-    let mut keys: Vec<(ChunkKey, i32)> = tree
+    let keys: Vec<IVec3Pod> = tree
         .occupied_chunks
         .iter()
-        .filter_map(|key| {
-            let dx = key.0 - centre.0;
-            let dy = key.1 - centre.1;
-            let dz = key.2 - centre.2;
-            if dx.abs() > r || dy.abs() > r || dz.abs() > r { return None; }
-            if spawned.0.contains_key(key) { return None; }
-            Some((*key, dx*dx + dy*dy + dz*dz))
-        })
+        .map(|k| IVec3Pod { x: k.0, y: k.1, z: k.2, _pad: 0 })
         .collect();
 
-    keys.sort_by_key(|(_, d)| *d);
+    let count = keys.len().min(MAX_VISIBLE_CHUNKS);
+    worker.write_slice("keys_in", &keys[..count]);
 
-    queue.keys.clear();
-    queue.set.clear();
-    for (key, _) in keys {
-        queue.keys.push_back(key);
-        queue.set.insert(key);
-    }
+    let params = VisibleParams {
+        centre: IVec3Pod { x: centre.0, y: centre.1, z: centre.2, _pad: 0 },
+        radius: cfg.view_distance_chunks,
+        count: count as u32,
+        _pad0: 0,
+    };
+    worker.write("params", &params);
+    worker.execute();
+    count_res.0 = count;
 }
 
 /// move a limited number of keys from the queue into the octree’s dirty set
@@ -61,4 +61,35 @@ pub fn process_chunk_queue(
             tree.dirty_chunks.insert(key);
         } else { break; }
     }
+}
+
+pub fn apply_visible_chunk_results(
+    mut worker    : ResMut<AppComputeWorker<VisibleChunksWorker>>,
+    mut queue     : ResMut<ChunkQueue>,
+    spawned       : Res<SpawnedChunks>,
+    count_res     : Res<VisibleChunkCount>,
+) {
+    if !worker.ready() {
+        return;
+    }
+
+    let mut results: Vec<ChunkResult> = worker.read_vec("results");
+    results.truncate(count_res.0);
+    let mut keys: Vec<(ChunkKey, i32)> = results
+        .into_iter()
+        .filter_map(|r| {
+            if r.dist2 < 0 { return None; }
+            let key = (r.key.x, r.key.y, r.key.z);
+            if spawned.0.contains_key(&key) { return None; }
+            Some((key, r.dist2))
+        })
+        .collect();
+    keys.sort_by_key(|(_, d)| *d);
+
+    queue.keys.clear();
+    queue.set.clear();
+    for (key, _) in keys {
+        queue.keys.push_back(key);
+        queue.set.insert(key);
+    }
 }

client/src/plugins/environment/systems/voxels/visible_chunks_compute.rs

@@ -0,0 +1,64 @@
+use bevy::prelude::*;
+use bevy_easy_compute::prelude::*;
+use bytemuck::{Pod, Zeroable};
+
+pub const MAX_VISIBLE_CHUNKS: usize = 4096;
+
+#[repr(C)]
+#[derive(Clone, Copy, Pod, Zeroable, ShaderType)]
+pub struct IVec3Pod {
+    pub x: i32,
+    pub y: i32,
+    pub z: i32,
+    pub _pad: i32,
+}
+
+#[repr(C)]
+#[derive(Clone, Copy, Pod, Zeroable, ShaderType)]
+pub struct VisibleParams {
+    pub centre: IVec3Pod,
+    pub radius: i32,
+    pub count: u32,
+    pub _pad0: u32,
+}
+
+#[repr(C)]
+#[derive(Clone, Copy, Pod, Zeroable)]
+pub struct ChunkResult {
+    pub key: IVec3Pod,
+    pub dist2: i32,
+}
+
+#[derive(TypePath)]
+struct VisibleShader;
+
+impl ComputeShader for VisibleShader {
+    fn shader() -> ShaderRef {
+        "shaders/visible_chunks.wgsl".into()
+    }
+}
+
+#[derive(Resource)]
+pub struct VisibleChunksWorker;
+
+impl ComputeWorker for VisibleChunksWorker {
+    fn build(world: &mut World) -> AppComputeWorker<Self> {
+        let default_params = VisibleParams {
+            centre: IVec3Pod { x: 0, y: 0, z: 0, _pad: 0 },
+            radius: 0,
+            count: 0,
+            _pad0: 0,
+        };
+        AppComputeWorkerBuilder::new(world)
+            .add_uniform("params", &default_params)
+            .add_staging("keys_in", &[IVec3Pod { x: 0, y: 0, z: 0, _pad: 0 }; MAX_VISIBLE_CHUNKS])
+            .add_staging("results", &[ChunkResult { key: IVec3Pod { x: 0, y: 0, z: 0, _pad: 0 }, dist2: 0 }; MAX_VISIBLE_CHUNKS])
+            .add_pass::<VisibleShader>([((MAX_VISIBLE_CHUNKS as u32 + 63) / 64), 1, 1], &["params", "keys_in", "results"])
+            .one_shot()
+            .synchronous()
+            .build()
+    }
+}
+
+#[derive(Resource, Default)]
+pub struct VisibleChunkCount(pub usize);