From db4fc07f4962106e36707dcf9634ac9ff226f1bd Mon Sep 17 00:00:00 2001 From: Elias Stepanik <40958815+eliasstepanik@users.noreply.github.com> Date: Wed, 11 Jun 2025 12:13:19 +0200 Subject: [PATCH] Add compute shader for voxel generation --- client/Cargo.toml | 2 + client/assets/shaders/noise.wgsl | 42 +++++++++ client/assets/shaders/sphere.wgsl | 24 +++++ client/src/app.rs | 4 + .../plugins/environment/environment_plugin.rs | 10 ++- .../environment/systems/voxel_system.rs | 2 +- .../plugins/environment/systems/voxels/mod.rs | 2 + .../systems/voxels/noise_compute.rs | 41 +++++++++ .../systems/voxels/sphere_compute.rs | 90 +++++++++++++++++++ 9 files changed, 213 insertions(+), 4 deletions(-) create mode 100644 client/assets/shaders/noise.wgsl create mode 100644 client/assets/shaders/sphere.wgsl create mode 100644 client/src/plugins/environment/systems/voxels/noise_compute.rs create mode 100644 client/src/plugins/environment/systems/voxels/sphere_compute.rs diff --git a/client/Cargo.toml b/client/Cargo.toml index a29dc88..68f1e97 100644 --- a/client/Cargo.toml +++ b/client/Cargo.toml @@ -21,4 +21,6 @@ smallvec = "1.14.0" once_cell = "1.21.3" rayon = "1.10.0" bincode = "1.3" +bevy_easy_compute = "0.15" +bytemuck = { version = "1", features = ["derive"] } diff --git a/client/assets/shaders/noise.wgsl b/client/assets/shaders/noise.wgsl new file mode 100644 index 0000000..bfe9e1c --- /dev/null +++ b/client/assets/shaders/noise.wgsl @@ -0,0 +1,42 @@ +struct Params { + frequency: f32, + amplitude: f32, + width: u32, + depth: u32, +}; + +@group(0) @binding(0) +var params: Params; + +@group(0) @binding(1) +var heights: array; + +fn hash(p: vec2) -> f32 { + let dot_val = f32(p.x * 1271 + p.y * 3117); + return fract(sin(dot_val) * 43758.5453); +} + +fn noise(p: vec2) -> f32 { + let i = vec2(floor(p)); + let f = fract(p); + + let a = hash(i); + let b = hash(i + vec2(1,0)); + let c = hash(i + vec2(0,1)); + let d = hash(i + vec2(1,1)); + + let u = f * f * (3.0 - 2.0 * f); + + return mix(a, b, u.x) + (c - a) * u.y * (1.0 - u.x) + (d - b) * u.x * u.y; +} + +@compute @workgroup_size(8,8,1) +fn main(@builtin(global_invocation_id) id: vec3) { + if (id.x >= params.width || id.y >= params.depth) { + return; + } + let index = id.y * params.width + id.x; + let pos = vec2(f32(id.x), f32(id.y)) * params.frequency; + let n = noise(pos); + heights[index] = n * params.amplitude; +} diff --git a/client/assets/shaders/sphere.wgsl b/client/assets/shaders/sphere.wgsl new file mode 100644 index 0000000..fdcd64e --- /dev/null +++ b/client/assets/shaders/sphere.wgsl @@ -0,0 +1,24 @@ +struct Params { + radius: u32, + diameter: u32, +}; + +@group(0) @binding(0) +var params: Params; + +@group(0) @binding(1) +var voxels: array; + +@compute @workgroup_size(8,8,8) +fn main(@builtin(global_invocation_id) id: vec3) { + if (id.x >= params.diameter || id.y >= params.diameter || id.z >= params.diameter) { + return; + } + let r = f32(params.radius); + let cx = f32(id.x) - r; + let cy = f32(id.y) - r; + let cz = f32(id.z) - r; + let inside = select(0u, 1u, cx * cx + cy * cy + cz * cz <= r * r); + let index = id.z * params.diameter * params.diameter + id.y * params.diameter + id.x; + voxels[index] = inside; +} diff --git a/client/src/app.rs b/client/src/app.rs index 958713e..228fdf6 100644 --- a/client/src/app.rs +++ b/client/src/app.rs @@ -1,5 +1,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 bevy::prelude::*; pub struct AppPlugin; @@ -8,6 +10,8 @@ impl Plugin for AppPlugin { app.add_plugins(crate::plugins::ui::ui_plugin::UiPlugin); app.add_plugins(crate::plugins::big_space::big_space_plugin::BigSpaceIntegrationPlugin); app.add_plugins(crate::plugins::environment::environment_plugin::EnvironmentPlugin); + app.add_plugins(AppComputePlugin); + app.add_plugins(AppComputeWorkerPlugin::::default()); //app.add_plugins(crate::plugins::network::network_plugin::NetworkPlugin); app.add_plugins(crate::plugins::input::input_plugin::InputPlugin); app.add_plugins(WireframePlugin); diff --git a/client/src/plugins/environment/environment_plugin.rs b/client/src/plugins/environment/environment_plugin.rs index f6ae4e7..e1080ca 100644 --- a/client/src/plugins/environment/environment_plugin.rs +++ b/client/src/plugins/environment/environment_plugin.rs @@ -1,5 +1,7 @@ use bevy::app::{App, Plugin, PreStartup, PreUpdate, Startup}; use bevy::prelude::*; +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}; @@ -16,8 +18,8 @@ impl Plugin for EnvironmentPlugin { ( crate::plugins::environment::systems::camera_system::setup, crate::plugins::environment::systems::environment_system::setup.after(crate::plugins::environment::systems::camera_system::setup), - crate::plugins::environment::systems::voxel_system::setup - + crate::plugins::environment::systems::voxel_system::setup, + execute_sphere_once.after(crate::plugins::environment::systems::voxel_system::setup), ), ); @@ -25,6 +27,7 @@ impl Plugin for EnvironmentPlugin { app.insert_resource(ChunkCullingCfg { view_distance_chunks }); app.insert_resource(ChunkBudget { per_frame: 20 }); app.init_resource::(); + app.init_resource::(); app.add_systems(Update, log_mesh_count); app // ------------------------------------------------------------------------ @@ -42,7 +45,8 @@ impl Plugin for EnvironmentPlugin { enqueue_visible_chunks, process_chunk_queue.after(enqueue_visible_chunks), update_chunk_lods.after(process_chunk_queue), - rebuild_dirty_chunks .after(process_chunk_queue), // 4. (re)mesh dirty chunks + rebuild_dirty_chunks .after(process_chunk_queue), + apply_sphere_result.after(rebuild_dirty_chunks), // 4. (re)mesh dirty chunks /* ---------- optional debug drawing ------- */ visualize_octree_system diff --git a/client/src/plugins/environment/systems/voxel_system.rs b/client/src/plugins/environment/systems/voxel_system.rs index 72ed8da..c3e38eb 100644 --- a/client/src/plugins/environment/systems/voxel_system.rs +++ b/client/src/plugins/environment/systems/voxel_system.rs @@ -47,7 +47,7 @@ pub fn setup( generate_voxel_sphere_parallel(&mut tree, center, radius, color); }*/ - generate_voxel_sphere(&mut tree, 200, color); + // voxel sphere is generated asynchronously via compute shader tree }; diff --git a/client/src/plugins/environment/systems/voxels/mod.rs b/client/src/plugins/environment/systems/voxels/mod.rs index 61b64e4..b6a7c20 100644 --- a/client/src/plugins/environment/systems/voxels/mod.rs +++ b/client/src/plugins/environment/systems/voxels/mod.rs @@ -9,3 +9,5 @@ pub mod render_chunks; pub mod culling; pub mod queue_systems; pub mod lod; +pub mod noise_compute; +pub mod sphere_compute; diff --git a/client/src/plugins/environment/systems/voxels/noise_compute.rs b/client/src/plugins/environment/systems/voxels/noise_compute.rs new file mode 100644 index 0000000..4b3f570 --- /dev/null +++ b/client/src/plugins/environment/systems/voxels/noise_compute.rs @@ -0,0 +1,41 @@ +use bevy::prelude::*; +use bevy_easy_compute::prelude::*; + +#[derive(ShaderType, Clone, Copy)] +pub struct NoiseParams { + pub frequency: f32, + pub amplitude: f32, + pub width: u32, + pub depth: u32, +} + +#[derive(TypePath)] +struct NoiseShader; + +impl ComputeShader for NoiseShader { + fn shader() -> ShaderRef { + "shaders/noise.wgsl".into() + } +} + +#[derive(Resource)] +pub struct NoiseWorker; + +impl ComputeWorker for NoiseWorker { + fn build(world: &mut World) -> AppComputeWorker { + let params = NoiseParams { + frequency: 0.1, + amplitude: 1.0, + width: 0, + depth: 0, + }; + + AppComputeWorkerBuilder::new(world) + .add_uniform("params", ¶ms) + .add_staging("heights", &[0.0_f32; 1]) + .add_pass::([1, 1, 1], &["params", "heights"]) + .one_shot() + .build() + } +} + diff --git a/client/src/plugins/environment/systems/voxels/sphere_compute.rs b/client/src/plugins/environment/systems/voxels/sphere_compute.rs new file mode 100644 index 0000000..10676cd --- /dev/null +++ b/client/src/plugins/environment/systems/voxels/sphere_compute.rs @@ -0,0 +1,90 @@ +use bevy::prelude::*; +use bevy_easy_compute::prelude::*; +use bytemuck::{Pod, Zeroable}; +use super::structure::{SparseVoxelOctree, Voxel}; + +#[repr(C)] +#[derive(ShaderType, Clone, Copy, Pod, Zeroable)] +pub struct SphereParams { + pub radius: u32, + pub diameter: u32, +} + +#[derive(TypePath)] +struct SphereShader; + +impl ComputeShader for SphereShader { + fn shader() -> ShaderRef { + "shaders/sphere.wgsl".into() + } +} + +#[derive(Resource)] +pub struct SphereWorker; + +impl ComputeWorker for SphereWorker { + fn build(world: &mut World) -> AppComputeWorker { + let radius = 200u32; + let diameter = radius * 2 + 1; + let params = SphereParams { radius, diameter }; + let buffer = vec![0u32; (diameter * diameter * diameter) as usize]; + + AppComputeWorkerBuilder::new(world) + .add_uniform("params", ¶ms) + .add_staging("voxels", &buffer) + .add_pass::([ + (diameter + 7) / 8, + (diameter + 7) / 8, + (diameter + 7) / 8, + ], &["params", "voxels"]) + .synchronous() + .one_shot() + .build() + } +} + +#[derive(Resource, Default)] +pub struct SphereGenerated(pub bool); + +pub fn execute_sphere_once( + mut worker: ResMut>, + mut generated: ResMut, +) { + if generated.0 { + return; + } + worker.execute(); + generated.0 = true; +} + +pub fn apply_sphere_result( + mut worker: ResMut>, + mut octree_q: Query<&mut SparseVoxelOctree>, + mut generated: ResMut, +) { + if !generated.0 || !worker.ready() { + return; + } + + let params: SphereParams = worker.read("params"); + let voxels: Vec = worker.read_vec("voxels"); + let mut octree = octree_q.single_mut(); + let step = octree.get_spacing_at_depth(octree.max_depth); + let radius = params.radius as i32; + let diameter = params.diameter as i32; + for x in 0..diameter { + for y in 0..diameter { + for z in 0..diameter { + let idx = (z * diameter * diameter + y * diameter + x) as usize; + if voxels[idx] != 0 { + let wx = (x - radius) as f32 * step; + let wy = (y - radius) as f32 * step; + let wz = (z - radius) as f32 * step; + octree.insert(Vec3::new(wx, wy, wz), Voxel { color: Color::rgb(0.2, 0.8, 0.2) }); + } + } + } + } + generated.0 = false; +} +