Add GPU meshing pipeline skeleton

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<f32>,
+    step: f32,
+    axis: u32,
+    dir: i32,
+    slice: u32,
+    n: vec3<f32>,
+};
+
+struct Vertex {
+    pos: vec3<f32>,
+    normal: vec3<f32>,
+    uv: vec2<f32>,
+};
+
+@group(0) @binding(0) var<storage, read> voxels: array<u32>;
+@group(0) @binding(1) var<uniform> params: Params;
+@group(0) @binding(2) var<storage, read_write> vertices: array<Vertex>;
+@group(0) @binding(3) var<storage, read_write> indices: array<u32>;
+@group(0) @binding(4) var<storage, read_write> counts: atomic<u32>;
+
+const N: u32 = 16u;
+
+@compute @workgroup_size(1)
+fn main(@builtin(global_invocation_id) id: vec3<u32>) {
+    // TODO: implement full greedy algorithm.
+    // This shader currently only reserves space for CPU-driven meshing.
+}

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 {

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::<GpuMeshingPipeline>()
+            .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::<AssetServer>();
+        let shader: Handle<Shader> = asset_server.load("shaders/greedy_meshing.wgsl");
+        let render_device = world.resource::<RenderDevice>();
+        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::<RenderDevice>();
+        let pipeline_cache = world.resource::<PipelineCache>();
+        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<MeshBufferPool>,
+    _pipeline: Res<GpuMeshingPipeline>,
+) {
+    // TODO: upload voxel buffers and dispatch compute passes per chunk.
+}

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;