From 1802595f7edd12d42f4e891eba3aa438e1a9afcc Mon Sep 17 00:00:00 2001 From: Elias Stepanik <40958815+eliasstepanik@users.noreply.github.com> Date: Fri, 13 Jun 2025 03:14:02 +0200 Subject: [PATCH] Add GPU meshing pipeline skeleton --- client/assets/shaders/greedy_meshing.wgsl | 32 ++++++ .../plugins/environment/environment_plugin.rs | 2 + .../environment/systems/voxels/meshing_gpu.rs | 107 ++++++++++++++++++ .../plugins/environment/systems/voxels/mod.rs | 7 +- 4 files changed, 145 insertions(+), 3 deletions(-) create mode 100644 client/assets/shaders/greedy_meshing.wgsl create mode 100644 client/src/plugins/environment/systems/voxels/meshing_gpu.rs diff --git a/client/assets/shaders/greedy_meshing.wgsl b/client/assets/shaders/greedy_meshing.wgsl new file mode 100644 index 0000000..e265ad8 --- /dev/null +++ b/client/assets/shaders/greedy_meshing.wgsl @@ -0,0 +1,32 @@ +// Generates mesh quads for a voxel chunk using a simple greedy algorithm. +// Each invocation processes a slice of the chunk along one axis. +// Results are stored in a vertex/index buffer. + +struct Params { + origin: vec3, + step: f32, + axis: u32, + dir: i32, + slice: u32, + n: vec3, +}; + +struct Vertex { + pos: vec3, + normal: vec3, + uv: vec2, +}; + +@group(0) @binding(0) var voxels: array; +@group(0) @binding(1) var params: Params; +@group(0) @binding(2) var vertices: array; +@group(0) @binding(3) var indices: array; +@group(0) @binding(4) var counts: atomic; + +const N: u32 = 16u; + +@compute @workgroup_size(1) +fn main(@builtin(global_invocation_id) id: vec3) { + // TODO: implement full greedy algorithm. + // This shader currently only reserves space for CPU-driven meshing. +} diff --git a/client/src/plugins/environment/environment_plugin.rs b/client/src/plugins/environment/environment_plugin.rs index be6e039..e78a1b4 100644 --- a/client/src/plugins/environment/environment_plugin.rs +++ b/client/src/plugins/environment/environment_plugin.rs @@ -1,5 +1,6 @@ use crate::plugins::environment::systems::voxels::debug::{draw_grid, visualize_octree_system}; use crate::plugins::environment::systems::voxels::lod::update_chunk_lods; +use crate::plugins::environment::systems::voxels::meshing_gpu::GpuMeshingPlugin; use crate::plugins::environment::systems::voxels::queue_systems; use crate::plugins::environment::systems::voxels::queue_systems::{ enqueue_visible_chunks, process_chunk_queue, @@ -24,6 +25,7 @@ impl Plugin for EnvironmentPlugin { crate::plugins::environment::systems::voxel_system::setup, ), ); + app.add_plugins(GpuMeshingPlugin); let view_distance_chunks = 100; app.insert_resource(ChunkCullingCfg { diff --git a/client/src/plugins/environment/systems/voxels/meshing_gpu.rs b/client/src/plugins/environment/systems/voxels/meshing_gpu.rs new file mode 100644 index 0000000..c0c2f65 --- /dev/null +++ b/client/src/plugins/environment/systems/voxels/meshing_gpu.rs @@ -0,0 +1,107 @@ +use bevy::prelude::*; +use bevy::render::render_resource::*; +use bevy::render::renderer::RenderDevice; +use bevy::render::RenderApp; + +use super::structure::{MeshBufferPool, SparseVoxelOctree}; + +/// Runs greedy meshing on the GPU. +pub struct GpuMeshingPlugin; + +impl Plugin for GpuMeshingPlugin { + fn build(&self, app: &mut App) { + let render_app = app.sub_app_mut(RenderApp); + render_app + .init_resource::() + .add_systems(Render, queue_gpu_meshing); + } +} + +#[derive(Resource)] +pub struct GpuMeshingPipeline { + pub pipeline: CachedComputePipelineId, + pub layout: BindGroupLayout, +} + +impl FromWorld for GpuMeshingPipeline { + fn from_world(world: &mut World) -> Self { + let asset_server = world.resource::(); + let shader: Handle = asset_server.load("shaders/greedy_meshing.wgsl"); + let render_device = world.resource::(); + let layout = render_device.create_bind_group_layout(&BindGroupLayoutDescriptor { + label: Some("meshing_layout"), + entries: &[ + BindGroupLayoutEntry { + binding: 0, + visibility: ShaderStages::COMPUTE, + ty: BindingType::Buffer { + ty: BufferBindingType::Storage { read_only: true }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + BindGroupLayoutEntry { + binding: 1, + visibility: ShaderStages::COMPUTE, + ty: BindingType::Buffer { + ty: BufferBindingType::Uniform, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + BindGroupLayoutEntry { + binding: 2, + visibility: ShaderStages::COMPUTE, + ty: BindingType::Buffer { + ty: BufferBindingType::Storage { read_only: false }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + BindGroupLayoutEntry { + binding: 3, + visibility: ShaderStages::COMPUTE, + ty: BindingType::Buffer { + ty: BufferBindingType::Storage { read_only: false }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + BindGroupLayoutEntry { + binding: 4, + visibility: ShaderStages::COMPUTE, + ty: BindingType::Buffer { + ty: BufferBindingType::Storage { read_only: false }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + ], + }); + let pipeline_descriptor = ComputePipelineDescriptor { + label: Some("meshing_pipeline".into()), + layout: vec![layout.clone()], + shader, + shader_defs: vec![], + entry_point: "main".into(), + }; + let render_queue = world.resource::(); + let pipeline_cache = world.resource::(); + let pipeline = pipeline_cache.queue_compute_pipeline(render_queue, &pipeline_descriptor); + Self { pipeline, layout } + } +} + +/// System that dispatches the compute shader for dirty chunks. +fn queue_gpu_meshing( + _octrees: Query<&SparseVoxelOctree>, + _pool: ResMut, + _pipeline: Res, +) { + // TODO: upload voxel buffers and dispatch compute passes per chunk. +} diff --git a/client/src/plugins/environment/systems/voxels/mod.rs b/client/src/plugins/environment/systems/voxels/mod.rs index 61b64e4..0b05e7c 100644 --- a/client/src/plugins/environment/systems/voxels/mod.rs +++ b/client/src/plugins/environment/systems/voxels/mod.rs @@ -4,8 +4,9 @@ pub mod octree; pub mod structure; mod chunk; -mod meshing; -pub mod render_chunks; pub mod culling; -pub mod queue_systems; pub mod lod; +mod meshing; +mod meshing_gpu; +pub mod queue_systems; +pub mod render_chunks;