From 45892040979d8ad5a53fe2b3c60e6eb3f298b728 Mon Sep 17 00:00:00 2001 From: Elias Stepanik <40958815+eliasstepanik@users.noreply.github.com> Date: Fri, 13 Jun 2025 12:45:29 +0200 Subject: [PATCH] Add GPU-based visible chunk culling --- client/assets/shaders/chunk_visibility.wgsl | 29 +++++++ .../plugins/environment/environment_plugin.rs | 10 ++- .../plugins/environment/systems/voxels/mod.rs | 1 + .../systems/voxels/visibility_gpu.rs | 85 +++++++++++++++++++ 4 files changed, 121 insertions(+), 4 deletions(-) create mode 100644 client/assets/shaders/chunk_visibility.wgsl create mode 100644 client/src/plugins/environment/systems/voxels/visibility_gpu.rs diff --git a/client/assets/shaders/chunk_visibility.wgsl b/client/assets/shaders/chunk_visibility.wgsl new file mode 100644 index 0000000..e25c196 --- /dev/null +++ b/client/assets/shaders/chunk_visibility.wgsl @@ -0,0 +1,29 @@ +// Computes visible chunk keys based on camera centre and view radius. +// Input arrays must match in length and are processed per invocation. + +struct Params { + centre: vec3; + radius: i32; + count: u32; +}; + +@group(0) @binding(0) var occupied: array>; +@group(0) @binding(1) var spawned: array; +@group(0) @binding(2) var out_keys: array>; +@group(0) @binding(3) var out_count: atomic; +@group(0) @binding(4) var params: Params; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) id: vec3) { + let idx = id.x; + if idx >= params.count { return; } + let key = occupied[idx]; + if spawned[idx] != 0u { return; } + let dx = key.x - params.centre.x; + let dy = key.y - params.centre.y; + let dz = key.z - params.centre.z; + if dx*dx + dy*dy + dz*dz <= params.radius * params.radius { + let i = atomicAdd(&out_count, 1u); + out_keys[i] = key; + } +} diff --git a/client/src/plugins/environment/environment_plugin.rs b/client/src/plugins/environment/environment_plugin.rs index 0dff8bd..9b21478 100644 --- a/client/src/plugins/environment/environment_plugin.rs +++ b/client/src/plugins/environment/environment_plugin.rs @@ -5,8 +5,9 @@ use crate::plugins::environment::systems::voxels::meshing_gpu::{ }; use bevy_app_compute::prelude::{AppComputePlugin, AppComputeWorkerPlugin}; 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::process_chunk_queue; +use crate::plugins::environment::systems::voxels::visibility_gpu::{ + enqueue_visible_chunks_gpu, GpuVisibilityWorker, }; use crate::plugins::environment::systems::voxels::render_chunks::rebuild_dirty_chunks; use crate::plugins::environment::systems::voxels::structure::{ @@ -30,6 +31,7 @@ impl Plugin for EnvironmentPlugin { ); app.add_plugins(AppComputePlugin); app.add_plugins(AppComputeWorkerPlugin::::default()); + app.add_plugins(AppComputeWorkerPlugin::::default()); let view_distance_chunks = 100; app.insert_resource(ChunkCullingCfg { @@ -52,8 +54,8 @@ impl Plugin for EnvironmentPlugin { Update, ( /* ---------- culling & streaming ---------- */ - enqueue_visible_chunks, - process_chunk_queue.after(enqueue_visible_chunks), + enqueue_visible_chunks_gpu, + process_chunk_queue.after(enqueue_visible_chunks_gpu), update_chunk_lods.after(process_chunk_queue), rebuild_dirty_chunks.after(process_chunk_queue), // 4. (re)mesh dirty chunks queue_gpu_meshing.after(rebuild_dirty_chunks), diff --git a/client/src/plugins/environment/systems/voxels/mod.rs b/client/src/plugins/environment/systems/voxels/mod.rs index b71c048..3c45954 100644 --- a/client/src/plugins/environment/systems/voxels/mod.rs +++ b/client/src/plugins/environment/systems/voxels/mod.rs @@ -8,5 +8,6 @@ pub mod culling; pub mod lod; mod meshing; pub mod meshing_gpu; +pub mod visibility_gpu; pub mod queue_systems; pub mod render_chunks; diff --git a/client/src/plugins/environment/systems/voxels/visibility_gpu.rs b/client/src/plugins/environment/systems/voxels/visibility_gpu.rs new file mode 100644 index 0000000..61c6ab6 --- /dev/null +++ b/client/src/plugins/environment/systems/voxels/visibility_gpu.rs @@ -0,0 +1,85 @@ +use bevy::prelude::*; +use bevy_app_compute::prelude::*; + +use super::structure::{ChunkCullingCfg, ChunkQueue, PrevCameraChunk, SpawnedChunks, SparseVoxelOctree}; +use crate::plugins::environment::systems::voxels::helper::world_to_chunk; + +#[repr(C)] +#[derive(ShaderType, Copy, Clone, Default)] +pub struct Params { + pub centre: IVec3, + pub radius: i32, + pub count: u32, +} + +#[derive(TypePath)] +struct VisibilityShader; + +impl ComputeShader for VisibilityShader { + fn shader() -> ShaderRef { + "shaders/chunk_visibility.wgsl".into() + } +} + +#[derive(Resource)] +pub struct GpuVisibilityWorker; + +impl ComputeWorker for GpuVisibilityWorker { + fn build(world: &mut World) -> AppComputeWorker { + AppComputeWorkerBuilder::new(world) + .add_storage::<[IVec3; 1]>("occupied", &[IVec3::ZERO; 1]) + .add_storage::<[u32; 1]>("spawned", &[0u32; 1]) + .add_rw_storage::<[IVec3; 1]>("out_keys", &[IVec3::ZERO; 1]) + .add_rw_storage::("out_count", &0u32) + .add_uniform("params", &Params::default()) + .add_pass::([1, 1, 1], &["occupied", "spawned", "out_keys", "out_count", "params"]) + .one_shot() + .build() + } +} + +/// GPU-driven implementation of `enqueue_visible_chunks`. +pub fn enqueue_visible_chunks_gpu( + mut worker: ResMut>, + tree_q: Query<&SparseVoxelOctree>, + cam_q: Query<&GlobalTransform, With>, + spawned: Res, + mut prev_cam: ResMut, + cfg: Res, + mut queue: ResMut, +) { + let Ok(tree) = tree_q.get_single() else { return }; + let Ok(cam_tf) = cam_q.get_single() else { return }; + let cam_pos = cam_tf.translation(); + let centre = world_to_chunk(tree, cam_pos); + if prev_cam.0 == Some(centre) { return; } + prev_cam.0 = Some(centre); + + if !worker.ready() { return; } + + let occupied: Vec = tree.occupied_chunks.iter().copied().collect(); + let mut spawned_flags = Vec::with_capacity(occupied.len()); + for key in &occupied { + spawned_flags.push(if spawned.0.contains_key(key) { 1u32 } else { 0u32 }); + } + worker.write_slice("occupied", &occupied); + worker.write_slice("spawned", &spawned_flags); + worker.write_slice("out_keys", &vec![IVec3::ZERO; occupied.len()]); + worker.write("out_count", &0u32); + + let params = Params { centre, radius: cfg.view_distance_chunks, count: occupied.len() as u32 }; + worker.write("params", ¶ms); + + let workgroups = ((occupied.len() as f32) / 64.0).ceil() as u32; + worker.pass(0).dispatch([workgroups, 1, 1]); + worker.execute(); + + let count: u32 = worker.read("out_count"); + let keys: Vec = worker.read_vec("out_keys"); + queue.keys.clear(); + queue.set.clear(); + for key in keys.into_iter().take(count as usize) { + queue.keys.push_back(key.into()); + queue.set.insert(key.into()); + } +}