From 0cf98496ed982601f7a9b73ac0254136a9e4fc9c Mon Sep 17 00:00:00 2001
From: Elias Stepanik <40958815+eliasstepanik@users.noreply.github.com>
Date: Fri, 13 Jun 2025 02:53:19 +0200
Subject: [PATCH 1/6] Add mesh buffer pooling

---
 .../plugins/environment/environment_plugin.rs | 45 +++++++------
 .../environment/systems/voxels/meshing.rs     | 63 ++++++++++++-------
 .../systems/voxels/render_chunks.rs           | 50 ++++++++-------
 .../environment/systems/voxels/structure.rs   | 62 ++++++++++++------
 4 files changed, 137 insertions(+), 83 deletions(-)

diff --git a/client/src/plugins/environment/environment_plugin.rs b/client/src/plugins/environment/environment_plugin.rs
index c28841c..be6e039 100644
--- a/client/src/plugins/environment/environment_plugin.rs
+++ b/client/src/plugins/environment/environment_plugin.rs
@@ -1,28 +1,34 @@
+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::queue_systems;
+use crate::plugins::environment::systems::voxels::queue_systems::{
+    enqueue_visible_chunks, process_chunk_queue,
+};
+use crate::plugins::environment::systems::voxels::render_chunks::rebuild_dirty_chunks;
+use crate::plugins::environment::systems::voxels::structure::{
+    ChunkBudget, ChunkCullingCfg, ChunkQueue, MeshBufferPool, PrevCameraChunk, SparseVoxelOctree,
+    SpawnedChunks,
+};
 use bevy::app::{App, Plugin, PreStartup, PreUpdate, Startup};
 use bevy::prelude::*;
-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::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};
 
 pub struct EnvironmentPlugin;
 impl Plugin for EnvironmentPlugin {
     fn build(&self, app: &mut App) {
-
         app.add_systems(
             Startup,
             (
                 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::environment_system::setup
+                    .after(crate::plugins::environment::systems::camera_system::setup),
+                crate::plugins::environment::systems::voxel_system::setup,
             ),
         );
 
         let view_distance_chunks = 100;
-        app.insert_resource(ChunkCullingCfg { view_distance_chunks });
+        app.insert_resource(ChunkCullingCfg {
+            view_distance_chunks,
+        });
         app.insert_resource(ChunkBudget { per_frame: 20 });
         app.init_resource::<PrevCameraChunk>();
         app.add_systems(Update, log_mesh_count);
@@ -32,6 +38,7 @@ impl Plugin for EnvironmentPlugin {
             // ------------------------------------------------------------------------
             .init_resource::<ChunkQueue>()
             .init_resource::<SpawnedChunks>()
+            .init_resource::<MeshBufferPool>()
             // ------------------------------------------------------------------------
             // frame update
             // ------------------------------------------------------------------------
@@ -42,8 +49,7 @@ 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), // 4.  (re)mesh dirty chunks
                     /* ---------- optional debug drawing ------- */
                     visualize_octree_system
                         .run_if(should_visualize_octree)
@@ -52,9 +58,8 @@ impl Plugin for EnvironmentPlugin {
                         .run_if(should_draw_grid)
                         .after(visualize_octree_system),
                 )
-                    .chain(),     // make the whole tuple execute in this exact order
+                    .chain(), // make the whole tuple execute in this exact order
             );
-
     }
 }
 
@@ -65,11 +70,15 @@ fn log_mesh_count(meshes: Res<Assets<Mesh>>, time: Res<Time>) {
 }
 
 fn should_visualize_octree(octree_query: Query<&SparseVoxelOctree>) -> bool {
-    let Ok(octree) = octree_query.get_single() else { return false };
+    let Ok(octree) = octree_query.get_single() else {
+        return false;
+    };
     octree.show_wireframe
 }
 
 fn should_draw_grid(octree_query: Query<&SparseVoxelOctree>) -> bool {
-    let Ok(octree) = octree_query.get_single() else { return false };
+    let Ok(octree) = octree_query.get_single() else {
+        return false;
+    };
     octree.show_world_grid
-}
\ No newline at end of file
+}
diff --git a/client/src/plugins/environment/systems/voxels/meshing.rs b/client/src/plugins/environment/systems/voxels/meshing.rs
index c0c3e63..5452d0e 100644
--- a/client/src/plugins/environment/systems/voxels/meshing.rs
+++ b/client/src/plugins/environment/systems/voxels/meshing.rs
@@ -1,7 +1,7 @@
+use crate::plugins::environment::systems::voxels::structure::*;
 use bevy::asset::RenderAssetUsages;
 use bevy::prelude::*;
-use bevy::render::mesh::{Indices, PrimitiveTopology, VertexAttributeValues, Mesh};
-use crate::plugins::environment::systems::voxels::structure::*;
+use bevy::render::mesh::{Indices, Mesh, PrimitiveTopology, VertexAttributeValues};
 
 /*pub(crate) fn mesh_chunk(
     buffer: &[[[Option<Voxel>; CHUNK_SIZE as usize]; CHUNK_SIZE as usize]; CHUNK_SIZE as usize],
@@ -96,7 +96,7 @@ use crate::plugins::environment::systems::voxels::structure::*;
             }
         }
     }
-    
+
     // ------  2nd pass :  +Z faces  ---------------------------------------------
     for z in 0..CHUNK_SIZE {                         //   +Z faces (normal +Z)
         let nz          =  1;
@@ -298,12 +298,12 @@ use crate::plugins::environment::systems::voxels::structure::*;
     mesh
 }*/
 
-
 pub(crate) fn mesh_chunk(
     buffer: &[[[Option<Voxel>; CHUNK_SIZE as usize]; CHUNK_SIZE as usize]; CHUNK_SIZE as usize],
     origin: Vec3,
-    step:   f32,
-    tree:   &SparseVoxelOctree,
+    step: f32,
+    tree: &SparseVoxelOctree,
+    pool: &mut MeshBufferPool,
 ) -> Option<Mesh> {
     // ────────────────────────────────────────────────────────────────────────────
     // Helpers
@@ -328,12 +328,18 @@ pub(crate) fn mesh_chunk(
     // Push a single quad (4 vertices, 6 indices).  `base` is the lower‑left
     // corner in world space; `u`/`v` are the tangent vectors (length 1); `size`
     // is expressed in world units along those axes; `n` is the face normal.
-    // Preallocate vertex buffers for better performance
+    // Preallocate vertex buffers for better performance, reusing the pool.
+    pool.clear();
     let voxel_count = N * N * N;
-    let mut positions = Vec::<[f32; 3]>::with_capacity(voxel_count * 4);
-    let mut normals   = Vec::<[f32; 3]>::with_capacity(voxel_count * 4);
-    let mut uvs       = Vec::<[f32; 2]>::with_capacity(voxel_count * 4);
-    let mut indices   = Vec::<u32>::with_capacity(voxel_count * 6);
+    pool.positions.reserve(voxel_count * 4);
+    pool.normals.reserve(voxel_count * 4);
+    pool.uvs.reserve(voxel_count * 4);
+    pool.indices.reserve(voxel_count * 6);
+
+    let positions = &mut pool.positions;
+    let normals = &mut pool.normals;
+    let uvs = &mut pool.uvs;
+    let indices = &mut pool.indices;
 
     let mut push_quad = |base: Vec3, size: Vec2, n: Vec3, u: Vec3, v: Vec3| {
         let i0 = positions.len() as u32;
@@ -361,7 +367,7 @@ pub(crate) fn mesh_chunk(
 
     // Axes: 0→X, 1→Y, 2→Z.  For each axis we process the negative and positive
     // faces (dir = −1 / +1).
-    for (axis, dir) in [ (0, -1), (0, 1), (1, -1), (1, 1), (2, -1), (2, 1) ] {
+    for (axis, dir) in [(0, -1), (0, 1), (1, -1), (1, 1), (2, -1), (2, 1)] {
         // Mapping of (u,v) axes and their unit vectors in world space.
         let (u_axis, v_axis, face_normal, u_vec, v_vec) = match (axis, dir) {
             (0, d) => (1, 2, Vec3::new(d as f32, 0.0, 0.0), Vec3::Y, Vec3::Z),
@@ -386,15 +392,17 @@ pub(crate) fn mesh_chunk(
                     let mut cell = [0i32; 3];
                     let mut neighbor = [0i32; 3];
 
-                    cell [axis] = slice as i32 + if dir == 1 { -1 } else { 0 };
+                    cell[axis] = slice as i32 + if dir == 1 { -1 } else { 0 };
                     neighbor[axis] = cell[axis] + dir;
 
-                    cell [u_axis] = u as i32;
-                    cell [v_axis] = v as i32;
+                    cell[u_axis] = u as i32;
+                    cell[v_axis] = v as i32;
                     neighbor[u_axis] = u as i32;
                     neighbor[v_axis] = v as i32;
 
-                    if filled(cell[0], cell[1], cell[2]) && !filled(neighbor[0], neighbor[1], neighbor[2]) {
+                    if filled(cell[0], cell[1], cell[2])
+                        && !filled(neighbor[0], neighbor[1], neighbor[2])
+                    {
                         mask[idx(u, v)] = true;
                     }
                 }
@@ -409,7 +417,10 @@ pub(crate) fn mesh_chunk(
 
                     // Determine the rectangle width.
                     let mut width = 1;
-                    while u0 + width < N && mask[idx(u0 + width, v0)] && !visited[idx(u0 + width, v0)] {
+                    while u0 + width < N
+                        && mask[idx(u0 + width, v0)]
+                        && !visited[idx(u0 + width, v0)]
+                    {
                         width += 1;
                     }
 
@@ -417,7 +428,9 @@ pub(crate) fn mesh_chunk(
                     let mut height = 1;
                     'h: while v0 + height < N {
                         for du in 0..width {
-                            if !mask[idx(u0 + du, v0 + height)] || visited[idx(u0 + du, v0 + height)] {
+                            if !mask[idx(u0 + du, v0 + height)]
+                                || visited[idx(u0 + du, v0 + height)]
+                            {
                                 break 'h;
                             }
                         }
@@ -466,19 +479,23 @@ pub(crate) fn mesh_chunk(
         return None;
     }
 
-    let mut mesh = Mesh::new(PrimitiveTopology::TriangleList, RenderAssetUsages::default());
+    let mut mesh = Mesh::new(
+        PrimitiveTopology::TriangleList,
+        RenderAssetUsages::default(),
+    );
     mesh.insert_attribute(
         Mesh::ATTRIBUTE_POSITION,
-        VertexAttributeValues::Float32x3(positions),
+        VertexAttributeValues::Float32x3(positions.clone()),
     );
     mesh.insert_attribute(
         Mesh::ATTRIBUTE_NORMAL,
-        VertexAttributeValues::Float32x3(normals),
+        VertexAttributeValues::Float32x3(normals.clone()),
     );
     mesh.insert_attribute(
         Mesh::ATTRIBUTE_UV_0,
-        VertexAttributeValues::Float32x2(uvs),
+        VertexAttributeValues::Float32x2(uvs.clone()),
     );
-    mesh.insert_indices(Indices::U32(indices));
+    mesh.insert_indices(Indices::U32(indices.clone()));
+    pool.clear();
     Some(mesh)
 }
diff --git a/client/src/plugins/environment/systems/voxels/render_chunks.rs b/client/src/plugins/environment/systems/voxels/render_chunks.rs
index eb74e1d..993d8be 100644
--- a/client/src/plugins/environment/systems/voxels/render_chunks.rs
+++ b/client/src/plugins/environment/systems/voxels/render_chunks.rs
@@ -1,27 +1,30 @@
-use std::collections::HashMap;
-use std::fmt::format;
+use crate::plugins::big_space::big_space_plugin::RootGrid;
+use crate::plugins::environment::systems::voxels::meshing::mesh_chunk;
+use crate::plugins::environment::systems::voxels::structure::*;
 use bevy::pbr::wireframe::Wireframe;
 use bevy::prelude::*;
 use bevy::render::mesh::Mesh;
 use big_space::prelude::GridCell;
 use itertools::Itertools;
-use crate::plugins::big_space::big_space_plugin::RootGrid;
-use crate::plugins::environment::systems::voxels::meshing::mesh_chunk;
-use crate::plugins::environment::systems::voxels::structure::*;
+use std::collections::HashMap;
+use std::fmt::format;
 
 /// rebuilds meshes only for chunks flagged dirty by the octree
 pub fn rebuild_dirty_chunks(
-    mut commands : Commands,
-    mut octrees  : Query<&mut SparseVoxelOctree>,
-    mut meshes   : ResMut<Assets<Mesh>>,
+    mut commands: Commands,
+    mut octrees: Query<&mut SparseVoxelOctree>,
+    mut meshes: ResMut<Assets<Mesh>>,
     mut materials: ResMut<Assets<StandardMaterial>>,
-    chunk_q      : Query<(Entity,
-                          &Chunk,
-                          &Mesh3d,
-                          &MeshMaterial3d<StandardMaterial>,
-                          &ChunkLod)>,
-    mut spawned  : ResMut<SpawnedChunks>,
-    root         : Res<RootGrid>,
+    chunk_q: Query<(
+        Entity,
+        &Chunk,
+        &Mesh3d,
+        &MeshMaterial3d<StandardMaterial>,
+        &ChunkLod,
+    )>,
+    mut spawned: ResMut<SpawnedChunks>,
+    mut pool: ResMut<MeshBufferPool>,
+    root: Res<RootGrid>,
 ) {
     // map ChunkKey → (entity, mesh-handle, material-handle)
     let existing: HashMap<ChunkKey, (Entity, Handle<Mesh>, Handle<StandardMaterial>, u32)> =
@@ -39,8 +42,7 @@ pub fn rebuild_dirty_chunks(
         let mut bufs = Vec::new();
         for key in tree.dirty_chunks.iter().copied() {
             let lod = existing.get(&key).map(|v| v.3).unwrap_or(0);
-            let mut buf =
-                [[[None; CHUNK_SIZE as usize]; CHUNK_SIZE as usize]; CHUNK_SIZE as usize];
+            let mut buf = [[[None; CHUNK_SIZE as usize]; CHUNK_SIZE as usize]; CHUNK_SIZE as usize];
 
             let half = tree.size * 0.5;
             let step = tree.get_spacing_at_depth(tree.max_depth);
@@ -85,7 +87,7 @@ pub fn rebuild_dirty_chunks(
         for (key, buf, origin, step, lod) in bufs {
             if let Some((ent, mesh_h, _mat_h, _)) = existing.get(&key).cloned() {
                 // update mesh in-place; keeps old asset id
-                match mesh_chunk(&buf, origin, step, &tree) {
+                match mesh_chunk(&buf, origin, step, &tree, &mut pool) {
                     Some(new_mesh) => {
                         if let Some(mesh) = meshes.get_mut(&mesh_h) {
                             *mesh = new_mesh;
@@ -98,10 +100,10 @@ pub fn rebuild_dirty_chunks(
                         spawned.0.remove(&key);
                     }
                 }
-            } else if let Some(mesh) = mesh_chunk(&buf, origin, step, &tree) {
+            } else if let Some(mesh) = mesh_chunk(&buf, origin, step, &tree, &mut pool) {
                 // spawn brand-new chunk only if mesh has faces
                 let mesh_h = meshes.add(mesh);
-                let mat_h  = materials.add(StandardMaterial::default());
+                let mat_h = materials.add(StandardMaterial::default());
 
                 commands.entity(root.0).with_children(|p| {
                     let e = p
@@ -110,7 +112,11 @@ pub fn rebuild_dirty_chunks(
                             MeshMaterial3d(mat_h.clone()),
                             Transform::default(),
                             GridCell::ZERO,
-                            Chunk { key, voxels: Vec::new(), dirty: false },
+                            Chunk {
+                                key,
+                                voxels: Vec::new(),
+                                dirty: false,
+                            },
                             ChunkLod(lod),
                             /*Wireframe,*/
                         ))
@@ -122,4 +128,4 @@ pub fn rebuild_dirty_chunks(
 
         tree.clear_dirty_flags();
     }
-}
\ No newline at end of file
+}
diff --git a/client/src/plugins/environment/systems/voxels/structure.rs b/client/src/plugins/environment/systems/voxels/structure.rs
index 21457e1..b687d5a 100644
--- a/client/src/plugins/environment/systems/voxels/structure.rs
+++ b/client/src/plugins/environment/systems/voxels/structure.rs
@@ -1,7 +1,7 @@
-use std::collections::{HashMap, HashSet, VecDeque};
 use bevy::color::Color;
 use bevy::prelude::*;
-use serde::{Serialize, Deserialize, Serializer, Deserializer};
+use serde::{Deserialize, Deserializer, Serialize, Serializer};
+use std::collections::{HashMap, HashSet, VecDeque};
 
 fn serialize_color<S>(color: &Color, serializer: S) -> Result<S::Ok, S::Error>
 where
@@ -19,11 +19,13 @@ where
     Ok(Color::linear_rgba(arr[0], arr[1], arr[2], arr[3]))
 }
 
-
 /// Represents a single voxel with a color.
 #[derive(Debug, Clone, Copy, Component, PartialEq, Default, Serialize, Deserialize)]
 pub struct Voxel {
-    #[serde(serialize_with = "serialize_color", deserialize_with = "deserialize_color")]
+    #[serde(
+        serialize_with = "serialize_color",
+        deserialize_with = "deserialize_color"
+    )]
     pub color: Color,
 }
 
@@ -44,7 +46,6 @@ pub struct OctreeNode {
 /// Represents the root of the sparse voxel octree.
 #[derive(Debug, Component, Serialize, Deserialize, Clone)]
 pub struct SparseVoxelOctree {
-
     pub root: OctreeNode,
     pub max_depth: u32,
     pub size: f32,
@@ -77,13 +78,10 @@ impl OctreeNode {
 impl Voxel {
     /// Creates a new empty octree node.
     pub fn new(color: Color) -> Self {
-        Self {
-            color,
-        }
+        Self { color }
     }
 }
 
-
 pub const NEIGHBOR_OFFSETS: [(f32, f32, f32); 6] = [
     (-1.0, 0.0, 0.0), // Left
     (1.0, 0.0, 0.0),  // Right
@@ -93,7 +91,6 @@ pub const NEIGHBOR_OFFSETS: [(f32, f32, f32); 6] = [
     (0.0, 0.0, 1.0),  // Front
 ];
 
-
 #[derive(Debug)]
 pub struct Ray {
     pub origin: Vec3,
@@ -106,25 +103,22 @@ pub struct AABB {
     pub max: Vec3,
 }
 
-pub const CHUNK_SIZE: i32 = 16;         // 16×16×16 voxels
-pub const CHUNK_POW : u32 = 4;
+pub const CHUNK_SIZE: i32 = 16; // 16×16×16 voxels
+pub const CHUNK_POW: u32 = 4;
 
 #[derive(Component)]
 pub struct Chunk {
     pub key: ChunkKey,
-    pub voxels: Vec<(IVec3, Voxel)>,   // local coords 0‥15
+    pub voxels: Vec<(IVec3, Voxel)>, // local coords 0‥15
     pub dirty: bool,
-
 }
 
 #[derive(Component, Debug, Clone, Copy)]
 pub struct ChunkLod(pub u32);
 
-
 #[derive(Debug, Clone, Copy, Hash, PartialEq, Eq, Serialize, Deserialize)]
 pub struct ChunkKey(pub i32, pub i32, pub i32);
 
-
 /// maximum amount of *new* chunk meshes we are willing to create each frame
 #[derive(Resource)]
 pub struct ChunkBudget {
@@ -132,7 +126,7 @@ pub struct ChunkBudget {
 }
 impl Default for ChunkBudget {
     fn default() -> Self {
-        Self { per_frame: 4 }      // tweak to taste
+        Self { per_frame: 4 } // tweak to taste
     }
 }
 
@@ -140,7 +134,7 @@ impl Default for ChunkBudget {
 #[derive(Resource, Default)]
 pub struct ChunkQueue {
     pub keys: VecDeque<ChunkKey>,
-    pub set:  HashSet<ChunkKey>,
+    pub set: HashSet<ChunkKey>,
 }
 
 /// map “which chunk key already has an entity in the world?”
@@ -149,8 +143,16 @@ pub struct SpawnedChunks(pub HashMap<ChunkKey, Entity>);
 
 /// how big the cube around the player is, measured in chunks
 #[derive(Resource)]
-pub struct ChunkCullingCfg { pub view_distance_chunks: i32 }
-impl Default for ChunkCullingCfg { fn default() -> Self { Self { view_distance_chunks: 6 } } }
+pub struct ChunkCullingCfg {
+    pub view_distance_chunks: i32,
+}
+impl Default for ChunkCullingCfg {
+    fn default() -> Self {
+        Self {
+            view_distance_chunks: 6,
+        }
+    }
+}
 
 #[derive(Resource, Default)]
 pub struct PrevCameraChunk(pub Option<ChunkKey>);
@@ -172,3 +174,23 @@ impl ChunkOffsets {
         Self(offsets)
     }
 }
+
+/// Pool reused when constructing chunk meshes. Reusing the backing
+/// storage avoids frequent allocations when rebuilding many chunks.
+#[derive(Resource, Default)]
+pub struct MeshBufferPool {
+    pub positions: Vec<[f32; 3]>,
+    pub normals: Vec<[f32; 3]>,
+    pub uvs: Vec<[f32; 2]>,
+    pub indices: Vec<u32>,
+}
+
+impl MeshBufferPool {
+    /// Clears all buffers while keeping the allocated capacity.
+    pub fn clear(&mut self) {
+        self.positions.clear();
+        self.normals.clear();
+        self.uvs.clear();
+        self.indices.clear();
+    }
+}

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 2/6] 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<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.
+}
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::<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.
+}
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;

From 5a7269a4464558916624da7bc405ada695a436a3 Mon Sep 17 00:00:00 2001
From: Elias Stepanik <40958815+eliasstepanik@users.noreply.github.com>
Date: Fri, 13 Jun 2025 03:19:44 +0200
Subject: [PATCH 3/6] Integrate bevy_app_compute for GPU meshing

---
 client/Cargo.toml                             |   2 +
 .../plugins/environment/environment_plugin.rs |   9 +-
 .../environment/systems/voxels/meshing_gpu.rs | 140 +++++++-----------
 .../plugins/environment/systems/voxels/mod.rs |   2 +-
 4 files changed, 60 insertions(+), 93 deletions(-)

diff --git a/client/Cargo.toml b/client/Cargo.toml
index 57e4c92..a373ba7 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_app_compute = "0.16"
+bytemuck = { version = "1.14", features = ["derive"] }
 
diff --git a/client/src/plugins/environment/environment_plugin.rs b/client/src/plugins/environment/environment_plugin.rs
index e78a1b4..0dff8bd 100644
--- a/client/src/plugins/environment/environment_plugin.rs
+++ b/client/src/plugins/environment/environment_plugin.rs
@@ -1,6 +1,9 @@
 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::meshing_gpu::{
+    GpuMeshingWorker, queue_gpu_meshing,
+};
+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,
@@ -25,7 +28,8 @@ impl Plugin for EnvironmentPlugin {
                 crate::plugins::environment::systems::voxel_system::setup,
             ),
         );
-        app.add_plugins(GpuMeshingPlugin);
+        app.add_plugins(AppComputePlugin);
+        app.add_plugins(AppComputeWorkerPlugin::<GpuMeshingWorker>::default());
 
         let view_distance_chunks = 100;
         app.insert_resource(ChunkCullingCfg {
@@ -52,6 +56,7 @@ impl Plugin for EnvironmentPlugin {
                     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
+                    queue_gpu_meshing.after(rebuild_dirty_chunks),
                     /* ---------- optional debug drawing ------- */
                     visualize_octree_system
                         .run_if(should_visualize_octree)
diff --git a/client/src/plugins/environment/systems/voxels/meshing_gpu.rs b/client/src/plugins/environment/systems/voxels/meshing_gpu.rs
index c0c2f65..5435131 100644
--- a/client/src/plugins/environment/systems/voxels/meshing_gpu.rs
+++ b/client/src/plugins/environment/systems/voxels/meshing_gpu.rs
@@ -1,107 +1,67 @@
 use bevy::prelude::*;
-use bevy::render::render_resource::*;
-use bevy::render::renderer::RenderDevice;
-use bevy::render::RenderApp;
+use bevy_app_compute::prelude::*;
+use bytemuck::{Pod, Zeroable};
 
 use super::structure::{MeshBufferPool, SparseVoxelOctree};
 
-/// Runs greedy meshing on the GPU.
-pub struct GpuMeshingPlugin;
+#[repr(C)]
+#[derive(ShaderType, Copy, Clone, Pod, Zeroable, Default)]
+pub struct Params {
+    pub origin: Vec3,
+    pub step: f32,
+    pub axis: u32,
+    pub dir: i32,
+    pub slice: u32,
+    pub _pad: u32,
+}
 
-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);
+#[repr(C)]
+#[derive(ShaderType, Copy, Clone, Pod, Zeroable, Default)]
+pub struct VertexGpu {
+    pub pos: Vec3,
+    pub normal: Vec3,
+    pub uv: Vec2,
+}
+
+#[derive(TypePath)]
+struct GreedyMeshingShader;
+
+impl ComputeShader for GreedyMeshingShader {
+    fn shader() -> ShaderRef {
+        "shaders/greedy_meshing.wgsl".into()
     }
 }
 
+/// GPU worker executing greedy meshing for chunks.
 #[derive(Resource)]
-pub struct GpuMeshingPipeline {
-    pub pipeline: CachedComputePipelineId,
-    pub layout: BindGroupLayout,
-}
+pub struct GpuMeshingWorker;
 
-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 }
+impl ComputeWorker for GpuMeshingWorker {
+    fn build(world: &mut World) -> AppComputeWorker<Self> {
+        AppComputeWorkerBuilder::new(world)
+            .add_storage::<u32>("voxels", &[0u32; 1])
+            .add_uniform("params", &Params::default())
+            .add_storage::<VertexGpu>("vertices", &[VertexGpu::default(); 1])
+            .add_storage::<u32>("indices", &[0u32; 1])
+            .add_storage::<u32>("counts", &[0u32; 2])
+            .add_pass::<GreedyMeshingShader>(
+                [1, 1, 1],
+                &["voxels", "params", "vertices", "indices", "counts"],
+            )
+            .one_shot()
+            .build()
     }
 }
 
-/// System that dispatches the compute shader for dirty chunks.
-fn queue_gpu_meshing(
+/// Placeholder system that will dispatch the compute worker for dirty chunks.
+pub fn queue_gpu_meshing(
+    mut worker: ResMut<AppComputeWorker<GpuMeshingWorker>>,
     _octrees: Query<&SparseVoxelOctree>,
     _pool: ResMut<MeshBufferPool>,
-    _pipeline: Res<GpuMeshingPipeline>,
 ) {
-    // TODO: upload voxel buffers and dispatch compute passes per chunk.
+    if !worker.ready() {
+        return;
+    }
+    // TODO: populate the worker buffers with chunk data before dispatching.
+    worker.execute();
 }
diff --git a/client/src/plugins/environment/systems/voxels/mod.rs b/client/src/plugins/environment/systems/voxels/mod.rs
index 0b05e7c..b71c048 100644
--- a/client/src/plugins/environment/systems/voxels/mod.rs
+++ b/client/src/plugins/environment/systems/voxels/mod.rs
@@ -7,6 +7,6 @@ mod chunk;
 pub mod culling;
 pub mod lod;
 mod meshing;
-mod meshing_gpu;
+pub mod meshing_gpu;
 pub mod queue_systems;
 pub mod render_chunks;

From b355ac13a99d37dc61c90162336ac3de160d8186 Mon Sep 17 00:00:00 2001
From: Elias Stepanik <40958815+eliasstepanik@users.noreply.github.com>
Date: Fri, 13 Jun 2025 03:29:56 +0200
Subject: [PATCH 4/6] Fix GPU worker storage types

---
 .../src/plugins/environment/systems/voxels/meshing_gpu.rs | 8 ++++----
 1 file changed, 4 insertions(+), 4 deletions(-)

diff --git a/client/src/plugins/environment/systems/voxels/meshing_gpu.rs b/client/src/plugins/environment/systems/voxels/meshing_gpu.rs
index 5435131..51a4ad1 100644
--- a/client/src/plugins/environment/systems/voxels/meshing_gpu.rs
+++ b/client/src/plugins/environment/systems/voxels/meshing_gpu.rs
@@ -39,11 +39,11 @@ pub struct GpuMeshingWorker;
 impl ComputeWorker for GpuMeshingWorker {
     fn build(world: &mut World) -> AppComputeWorker<Self> {
         AppComputeWorkerBuilder::new(world)
-            .add_storage::<u32>("voxels", &[0u32; 1])
+            .add_storage("voxels", &[0u32; 1])
             .add_uniform("params", &Params::default())
-            .add_storage::<VertexGpu>("vertices", &[VertexGpu::default(); 1])
-            .add_storage::<u32>("indices", &[0u32; 1])
-            .add_storage::<u32>("counts", &[0u32; 2])
+            .add_storage("vertices", &[VertexGpu::default(); 1])
+            .add_storage("indices", &[0u32; 1])
+            .add_storage("counts", &[0u32; 2])
             .add_pass::<GreedyMeshingShader>(
                 [1, 1, 1],
                 &["voxels", "params", "vertices", "indices", "counts"],

From eb34a16a0adb10ca70e817ba450af2ad092f0450 Mon Sep 17 00:00:00 2001
From: Elias Stepanik <40958815+eliasstepanik@users.noreply.github.com>
Date: Fri, 13 Jun 2025 03:37:27 +0200
Subject: [PATCH 5/6] Fix GPU mesher structs

---
 client/src/plugins/environment/systems/voxels/meshing_gpu.rs | 5 ++---
 1 file changed, 2 insertions(+), 3 deletions(-)

diff --git a/client/src/plugins/environment/systems/voxels/meshing_gpu.rs b/client/src/plugins/environment/systems/voxels/meshing_gpu.rs
index 51a4ad1..13c60c5 100644
--- a/client/src/plugins/environment/systems/voxels/meshing_gpu.rs
+++ b/client/src/plugins/environment/systems/voxels/meshing_gpu.rs
@@ -1,11 +1,10 @@
 use bevy::prelude::*;
 use bevy_app_compute::prelude::*;
-use bytemuck::{Pod, Zeroable};
 
 use super::structure::{MeshBufferPool, SparseVoxelOctree};
 
 #[repr(C)]
-#[derive(ShaderType, Copy, Clone, Pod, Zeroable, Default)]
+#[derive(ShaderType, Copy, Clone, Default)]
 pub struct Params {
     pub origin: Vec3,
     pub step: f32,
@@ -16,7 +15,7 @@ pub struct Params {
 }
 
 #[repr(C)]
-#[derive(ShaderType, Copy, Clone, Pod, Zeroable, Default)]
+#[derive(ShaderType, Copy, Clone, Default)]
 pub struct VertexGpu {
     pub pos: Vec3,
     pub normal: Vec3,

From 6c4b125ea9bb44d6e8768c7be35ba1193b87d426 Mon Sep 17 00:00:00 2001
From: Elias Stepanik <40958815+eliasstepanik@users.noreply.github.com>
Date: Fri, 13 Jun 2025 12:24:23 +0200
Subject: [PATCH 6/6] Implement greedy meshing shader

---
 client/assets/shaders/greedy_meshing.wgsl | 144 +++++++++++++++++++++-
 1 file changed, 141 insertions(+), 3 deletions(-)

diff --git a/client/assets/shaders/greedy_meshing.wgsl b/client/assets/shaders/greedy_meshing.wgsl
index e265ad8..e13407f 100644
--- a/client/assets/shaders/greedy_meshing.wgsl
+++ b/client/assets/shaders/greedy_meshing.wgsl
@@ -8,7 +8,7 @@ struct Params {
     axis: u32,
     dir: i32,
     slice: u32,
-    n: vec3<f32>,
+    _pad: u32,
 };
 
 struct Vertex {
@@ -25,8 +25,146 @@ struct Vertex {
 
 const N: u32 = 16u;
 
+const MASK_LEN: u32 = N * N;
+
+fn voxel_index(p: vec3<i32>) -> u32 {
+    return u32(p.x) * N * N + u32(p.y) * N + u32(p.z);
+}
+
+fn voxel_filled(p: vec3<i32>) -> bool {
+    return p.x >= 0 && p.x < i32(N) && p.y >= 0 && p.y < i32(N) && p.z >= 0 && p.z < i32(N) && voxels[voxel_index(p)] != 0u;
+}
+
 @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.
+    var mask: array<bool, MASK_LEN>;
+    var visited: array<bool, MASK_LEN>;
+
+    // Iterate over all axes and both face directions.
+    for (var axis: u32 = 0u; axis < 3u; axis = axis + 1u) {
+        for (var dir_idx: u32 = 0u; dir_idx < 2u; dir_idx = dir_idx + 1u) {
+            let dir: i32 = select(-1, 1, dir_idx == 1u);
+
+            for (var slice: u32 = 0u; slice < N; slice = slice + 1u) {
+                // Build mask for this slice.
+                for (var u: u32 = 0u; u < N; u = u + 1u) {
+                    for (var v: u32 = 0u; v < N; v = v + 1u) {
+                        var cell = vec3<i32>(0, 0, 0);
+                        var neighbor = vec3<i32>(0, 0, 0);
+
+                        if axis == 0u {
+                            cell = vec3<i32>(i32(slice), i32(u), i32(v));
+                            neighbor = cell + vec3<i32>(dir, 0, 0);
+                        } else if axis == 1u {
+                            cell = vec3<i32>(i32(v), i32(slice), i32(u));
+                            neighbor = cell + vec3<i32>(0, dir, 0);
+                        } else {
+                            cell = vec3<i32>(i32(u), i32(v), i32(slice));
+                            neighbor = cell + vec3<i32>(0, 0, dir);
+                        }
+
+                        let i = u * N + v;
+                        mask[i] = voxel_filled(cell) && !voxel_filled(neighbor);
+                        visited[i] = false;
+                    }
+                }
+
+                // Greedy merge.
+                for (var u0: u32 = 0u; u0 < N; u0 = u0 + 1u) {
+                    for (var v0: u32 = 0u; v0 < N; v0 = v0 + 1u) {
+                        let i0 = u0 * N + v0;
+                        if !mask[i0] || visited[i0] {
+                            continue;
+                        }
+
+                        var width: u32 = 1u;
+                        loop {
+                            if u0 + width >= N || !mask[u0 + width * N + v0] || visited[u0 + width * N + v0] {
+                                break;
+                            }
+                            width = width + 1u;
+                        }
+
+                        var height: u32 = 1u;
+                        outer: loop {
+                            if v0 + height >= N {
+                                break;
+                            }
+                            for (var du: u32 = 0u; du < width; du = du + 1u) {
+                                let idx = (u0 + du) * N + v0 + height;
+                                if !mask[idx] || visited[idx] {
+                                    break outer;
+                                }
+                            }
+                            height = height + 1u;
+                        }
+
+                        for (var du: u32 = 0u; du < width; du = du + 1u) {
+                            for (var dv: u32 = 0u; dv < height; dv = dv + 1u) {
+                                visited[(u0 + du) * N + v0 + dv] = true;
+                            }
+                        }
+
+                        // Compute base world-space position.
+                        var base = params.origin;
+                        if axis == 0u {
+                            base = base + vec3<f32>(f32(slice) + (dir > 0 ? 1.0 : 0.0), f32(u0), f32(v0)) * params.step;
+                        } else if axis == 1u {
+                            base = base + vec3<f32>(f32(v0), f32(slice) + (dir > 0 ? 1.0 : 0.0), f32(u0)) * params.step;
+                        } else {
+                            base = base + vec3<f32>(f32(u0), f32(v0), f32(slice) + (dir > 0 ? 1.0 : 0.0)) * params.step;
+                        }
+
+                        let size = vec2<f32>(f32(width) * params.step, f32(height) * params.step);
+
+                        var normal = vec3<f32>(0.0, 0.0, 0.0);
+                        var u_unit = vec3<f32>(0.0, 0.0, 0.0);
+                        var v_unit = vec3<f32>(0.0, 0.0, 0.0);
+
+                        if axis == 0u {
+                            normal = vec3<f32>(f32(dir), 0.0, 0.0);
+                            u_unit = vec3<f32>(0.0, 1.0, 0.0);
+                            v_unit = vec3<f32>(0.0, 0.0, 1.0);
+                        } else if axis == 1u {
+                            normal = vec3<f32>(0.0, f32(dir), 0.0);
+                            u_unit = vec3<f32>(0.0, 0.0, 1.0);
+                            v_unit = vec3<f32>(1.0, 0.0, 0.0);
+                        } else {
+                            normal = vec3<f32>(0.0, 0.0, f32(dir));
+                            u_unit = vec3<f32>(1.0, 0.0, 0.0);
+                            v_unit = vec3<f32>(0.0, 1.0, 0.0);
+                        }
+
+                        let p0 = base;
+                        let p1 = base + u_unit * size.x;
+                        let p2 = base + u_unit * size.x + v_unit * size.y;
+                        let p3 = base + v_unit * size.y;
+
+                        let vi = atomicAdd(&counts[0], 4u);
+                        vertices[vi] = Vertex(pos: p0, normal: normal, uv: vec2<f32>(0.0, 1.0));
+                        vertices[vi + 1u] = Vertex(pos: p1, normal: normal, uv: vec2<f32>(1.0, 1.0));
+                        vertices[vi + 2u] = Vertex(pos: p2, normal: normal, uv: vec2<f32>(1.0, 0.0));
+                        vertices[vi + 3u] = Vertex(pos: p3, normal: normal, uv: vec2<f32>(0.0, 0.0));
+
+                        let ii = atomicAdd(&counts[1], 6u);
+                        if dir > 0 {
+                            indices[ii] = vi;
+                            indices[ii + 1u] = vi + 1u;
+                            indices[ii + 2u] = vi + 2u;
+                            indices[ii + 3u] = vi + 2u;
+                            indices[ii + 4u] = vi + 3u;
+                            indices[ii + 5u] = vi;
+                        } else {
+                            indices[ii] = vi;
+                            indices[ii + 1u] = vi + 3u;
+                            indices[ii + 2u] = vi + 2u;
+                            indices[ii + 3u] = vi + 2u;
+                            indices[ii + 4u] = vi + 1u;
+                            indices[ii + 5u] = vi;
+                        }
+                    }
+                }
+            }
+        }
+    }
 }