Add GPU meshing worker

This commit is contained in:
Elias Stepanik 2025-06-12 01:10:23 +02:00
parent 8534aaa6d7
commit 95c088cabb
5 changed files with 211 additions and 6 deletions

View File

@ -0,0 +1,119 @@
struct Params {
origin: vec3<f32>;
step: f32;
};
@group(0) @binding(0)
var<uniform> params: Params;
@group(0) @binding(1)
var<storage, read> voxels: array<u32>;
@group(0) @binding(2)
var<storage, read_write> positions: array<vec3<f32>>;
@group(0) @binding(3)
var<storage, read_write> normals: array<vec3<f32>>;
@group(0) @binding(4)
var<storage, read_write> uvs: array<vec2<f32>>;
@group(0) @binding(5)
var<storage, read_write> indices: array<u32>;
struct Counts {
vertex: atomic<u32>;
index: atomic<u32>;
};
@group(0) @binding(6)
var<storage, read_write> counts: Counts;
const N: u32 = 16u;
fn push_face(base: vec3<f32>, size: vec2<f32>, n: vec3<f32>, u: vec3<f32>, v: vec3<f32>) {
let vi = atomicAdd(&counts.vertex, 4u);
positions[vi + 0u] = base;
positions[vi + 1u] = base + u * size.x;
positions[vi + 2u] = base + u * size.x + v * size.y;
positions[vi + 3u] = base + v * size.y;
normals[vi + 0u] = n;
normals[vi + 1u] = n;
normals[vi + 2u] = n;
normals[vi + 3u] = n;
uvs[vi + 0u] = vec2<f32>(0.0, 1.0);
uvs[vi + 1u] = vec2<f32>(1.0, 1.0);
uvs[vi + 2u] = vec2<f32>(1.0, 0.0);
uvs[vi + 3u] = vec2<f32>(0.0, 0.0);
let ii = atomicAdd(&counts.index, 6u);
if (n.x + n.y + n.z >= 0.0) {
indices[ii + 0u] = 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 + 0u] = 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;
}
}
@compute @workgroup_size(1,1,1)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
if (gid.x > 0u) { return; }
let s = params.step;
for (var z: u32 = 0u; z < N; z = z + 1u) {
for (var y: u32 = 0u; y < N; y = y + 1u) {
for (var x: u32 = 0u; x < N; x = x + 1u) {
let idx = z * N * N + y * N + x;
if (voxels[idx] == 0u) { continue; }
let base = params.origin + vec3<f32>(f32(x) * s, f32(y) * s, f32(z) * s);
var filled: bool;
// -X
filled = false;
if (x > 0u) { filled = voxels[idx - 1u] != 0u; }
if (!filled) {
push_face(base, vec2<f32>(s, s), vec3<f32>(-1.0,0.0,0.0), vec3<f32>(0.0,0.0,s), vec3<f32>(0.0,s,0.0));
}
// +X
filled = false;
if (x + 1u < N) { filled = voxels[idx + 1u] != 0u; }
if (!filled) {
let b = base + vec3<f32>(s,0.0,0.0);
push_face(b, vec2<f32>(s, s), vec3<f32>(1.0,0.0,0.0), vec3<f32>(0.0,s,0.0), vec3<f32>(0.0,0.0,s));
}
// -Y
filled = false;
if (y > 0u) { filled = voxels[idx - N] != 0u; }
if (!filled) {
push_face(base, vec2<f32>(s, s), vec3<f32>(0.0,-1.0,0.0), vec3<f32>(s,0.0,0.0), vec3<f32>(0.0,0.0,s));
}
// +Y
filled = false;
if (y + 1u < N) { filled = voxels[idx + N] != 0u; }
if (!filled) {
let b = base + vec3<f32>(0.0,s,0.0);
push_face(b, vec2<f32>(s, s), vec3<f32>(0.0,1.0,0.0), vec3<f32>(0.0,0.0,s), vec3<f32>(s,0.0,0.0));
}
// -Z
filled = false;
if (z > 0u) { filled = voxels[idx - N * N] != 0u; }
if (!filled) {
push_face(base, vec2<f32>(s, s), vec3<f32>(0.0,0.0,-1.0), vec3<f32>(s,0.0,0.0), vec3<f32>(0.0,s,0.0));
}
// +Z
filled = false;
if (z + 1u < N) { filled = voxels[idx + N * N] != 0u; }
if (!filled) {
let b = base + vec3<f32>(0.0,0.0,s);
push_face(b, vec2<f32>(s, s), vec3<f32>(0.0,0.0,1.0), vec3<f32>(s,0.0,0.0), vec3<f32>(0.0,s,0.0));
}
}
}
}
}

View File

@ -3,6 +3,7 @@ use crate::helper::debug_gizmos::debug_gizmos;
use bevy_easy_compute::prelude::{AppComputePlugin, AppComputeWorkerPlugin};
use crate::plugins::environment::systems::voxels::sphere_compute::SphereWorker;
use crate::plugins::environment::systems::voxels::visible_chunks_compute::VisibleChunksWorker;
use crate::plugins::environment::systems::voxels::chunk_mesh_compute::ChunkMeshWorker;
use bevy::prelude::*;
pub struct AppPlugin;
@ -14,6 +15,7 @@ impl Plugin for AppPlugin {
app.add_plugins(AppComputePlugin);
app.add_plugins(AppComputeWorkerPlugin::<SphereWorker>::default());
app.add_plugins(AppComputeWorkerPlugin::<VisibleChunksWorker>::default());
app.add_plugins(AppComputeWorkerPlugin::<ChunkMeshWorker>::default());
//app.add_plugins(crate::plugins::network::network_plugin::NetworkPlugin);
app.add_plugins(crate::plugins::input::input_plugin::InputPlugin);
app.add_plugins(WireframePlugin);

View File

@ -0,0 +1,52 @@
use bevy::prelude::*;
use bevy_easy_compute::prelude::*;
use bytemuck::{Pod, Zeroable};
use super::structure::CHUNK_SIZE;
const MAX_FACES: usize = (CHUNK_SIZE as usize) * (CHUNK_SIZE as usize) * (CHUNK_SIZE as usize) * 6;
pub const MAX_VERTICES: usize = MAX_FACES * 4;
pub const MAX_INDICES: usize = MAX_FACES * 6;
#[repr(C)]
#[derive(Clone, Copy, Pod, Zeroable, ShaderType)]
pub struct MeshParams {
pub origin: [f32; 3],
pub step: f32,
}
#[repr(C)]
#[derive(Clone, Copy, Pod, Zeroable, ShaderType, Default)]
pub struct MeshCounts {
pub vertex_count: u32,
pub index_count: u32,
}
#[derive(TypePath)]
struct MeshShader;
impl ComputeShader for MeshShader {
fn shader() -> ShaderRef {
"shaders/chunk_mesh.wgsl".into()
}
}
#[derive(Resource)]
pub struct ChunkMeshWorker;
impl ComputeWorker for ChunkMeshWorker {
fn build(world: &mut World) -> AppComputeWorker<Self> {
let params = MeshParams { origin: [0.0; 3], step: 1.0 };
AppComputeWorkerBuilder::new(world)
.add_uniform("params", &params)
.add_staging("voxels", &[0u32; CHUNK_SIZE as usize * CHUNK_SIZE as usize * CHUNK_SIZE as usize])
.add_staging("positions", &[[0.0_f32; 3]; MAX_VERTICES])
.add_staging("normals", &[[0.0_f32; 3]; MAX_VERTICES])
.add_staging("uvs", &[[0.0_f32; 2]; MAX_VERTICES])
.add_staging("indices", &[0u32; MAX_INDICES])
.add_staging("counts", &MeshCounts::default())
.add_pass::<MeshShader>([1, 1, 1], &["params", "voxels", "positions", "normals", "uvs", "indices", "counts"])
.one_shot()
.synchronous()
.build()
}
}

View File

@ -12,3 +12,4 @@ pub mod lod;
pub mod noise_compute;
pub mod sphere_compute;
pub mod visible_chunks_compute;
pub mod chunk_mesh_compute;

View File

@ -1,12 +1,15 @@
use std::collections::HashMap;
use std::fmt::format;
use bevy::pbr::wireframe::Wireframe;
use bevy::prelude::*;
use bevy::render::mesh::Mesh;
use bevy::render::mesh::{
Mesh, PrimitiveTopology, Indices, VertexAttributeValues, RenderAssetUsages,
};
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::chunk_mesh_compute::{
ChunkMeshWorker, MeshParams, MeshCounts,
};
use bevy_easy_compute::prelude::AppComputeWorker;
use crate::plugins::environment::systems::voxels::structure::*;
/// rebuilds meshes only for chunks flagged dirty by the octree
@ -22,6 +25,7 @@ pub fn rebuild_dirty_chunks(
&ChunkLod)>,
mut spawned : ResMut<SpawnedChunks>,
root : Res<RootGrid>,
mut worker : ResMut<AppComputeWorker<ChunkMeshWorker>>,
) {
// map ChunkKey → (entity, mesh-handle, material-handle)
let existing: HashMap<ChunkKey, (Entity, Handle<Mesh>, Handle<StandardMaterial>, u32)> =
@ -83,9 +87,36 @@ pub fn rebuild_dirty_chunks(
//------------------------------------------------ create / update
for (key, buf, origin, step, lod) in bufs {
let voxels: Vec<u32> = buf
.iter()
.flat_map(|p| p.iter().flat_map(|r| r.iter().map(|v| if v.is_some() { 1u32 } else { 0u32 })))
.collect();
let params = MeshParams { origin: [origin.x, origin.y, origin.z], step };
worker.write("params", &params);
worker.write_slice("voxels", &voxels);
worker.write("counts", &MeshCounts::default());
worker.execute();
let counts: MeshCounts = worker.read("counts");
let mut positions: Vec<[f32; 3]> = worker.read_vec("positions");
let mut normals: Vec<[f32; 3]> = worker.read_vec("normals");
let mut uvs: Vec<[f32; 2]> = worker.read_vec("uvs");
let mut indices: Vec<u32> = worker.read_vec("indices");
positions.truncate(counts.vertex_count as usize);
normals.truncate(counts.vertex_count as usize);
uvs.truncate(counts.vertex_count as usize);
indices.truncate(counts.index_count as usize);
let mesh_opt = if counts.index_count > 0 {
let mut mesh = Mesh::new(PrimitiveTopology::TriangleList, RenderAssetUsages::default());
mesh.insert_attribute(Mesh::ATTRIBUTE_POSITION, VertexAttributeValues::Float32x3(positions));
mesh.insert_attribute(Mesh::ATTRIBUTE_NORMAL, VertexAttributeValues::Float32x3(normals));
mesh.insert_attribute(Mesh::ATTRIBUTE_UV_0, VertexAttributeValues::Float32x2(uvs));
mesh.insert_indices(Indices::U32(indices));
Some(mesh)
} else { None };
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_opt {
Some(new_mesh) => {
if let Some(mesh) = meshes.get_mut(&mesh_h) {
*mesh = new_mesh;
@ -98,7 +129,7 @@ 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_opt {
// spawn brand-new chunk only if mesh has faces
let mesh_h = meshes.add(mesh);
let mat_h = materials.add(StandardMaterial::default());