Meshlet continuous LOD (#12755)

Adds a basic level of detail system to meshlets. An extremely brief
summary is as follows:
* In `from_mesh.rs`, once we've built the first level of clusters, we
group clusters, simplify the new mega-clusters, and then split the
simplified groups back into regular sized clusters. Repeat several times
(ideally until you can't anymore). This forms a directed acyclic graph
(DAG), where the children are the meshlets from the previous level, and
the parents are the more simplified versions of their children. The leaf
nodes are meshlets formed from the original mesh.
* In `cull_meshlets.wgsl`, each cluster selects whether to render or not
based on the LOD bounding sphere (different than the culling bounding
sphere) of the current meshlet, the LOD bounding sphere of its parent
(the meshlet group from simplification), and the simplification error
relative to its children of both the current meshlet and its parent
meshlet. This kind of breaks two pass occlusion culling, which will be
fixed in a future PR by using an HZB from the previous frame to get the
initial list of occluders.

Many, _many_ improvements to be done in the future
https://github.com/bevyengine/bevy/issues/11518, not least of which is
code quality and speed. I don't even expect this to work on many types
of input meshes. This is just a basic implementation/draft for
collaboration.

Arguable how much we want to do in this PR, I'll leave that up to
maintainers. I've erred on the side of "as basic as possible".

References:
* Slides 27-77 (video available on youtube)
https://advances.realtimerendering.com/s2021/Karis_Nanite_SIGGRAPH_Advances_2021_final.pdf
*
https://blog.traverseresearch.nl/creating-a-directed-acyclic-graph-from-a-mesh-1329e57286e5
*
https://jglrxavpok.github.io/2024/01/19/recreating-nanite-lod-generation.html,
https://jglrxavpok.github.io/2024/03/12/recreating-nanite-faster-lod-generation.html,
https://jglrxavpok.github.io/2024/04/02/recreating-nanite-runtime-lod-selection.html,
and https://github.com/jglrxavpok/Carrot
*
https://github.com/gents83/INOX/tree/master/crates/plugins/binarizer/src
* https://cs418.cs.illinois.edu/website/text/nanite.html


![image](https://github.com/bevyengine/bevy/assets/47158642/e40bff9b-7d0c-4a19-a3cc-2aad24965977)

![image](https://github.com/bevyengine/bevy/assets/47158642/442c7da3-7761-4da7-9acd-37f15dd13e26)

---------

Co-authored-by: Ricky Taylor <rickytaylor26@gmail.com>
Co-authored-by: vero <email@atlasdostal.com>
Co-authored-by: François <mockersf@gmail.com>
Co-authored-by: atlas dostal <rodol@rivalrebels.com>
Co-authored-by: Patrick Walton <pcwalton@mimiga.net>
This commit is contained in:
JMS55 2024-04-23 14:43:53 -07:00 committed by GitHub
parent 17633c1f75
commit 6d6810c90d
No known key found for this signature in database
GPG key ID: B5690EEEBB952194
12 changed files with 491 additions and 159 deletions

Binary file not shown.

View file

@ -16,9 +16,15 @@ shader_format_glsl = ["bevy_render/shader_format_glsl"]
trace = ["bevy_render/trace"]
ios_simulator = ["bevy_render/ios_simulator"]
# Enables the meshlet renderer for dense high-poly scenes (experimental)
meshlet = ["dep:range-alloc", "dep:bincode"]
meshlet = [
"dep:lz4_flex",
"dep:serde",
"dep:bincode",
"dep:thiserror",
"dep:range-alloc",
]
# Enables processing meshes into meshlet meshes
meshlet_processor = ["dep:meshopt", "dep:thiserror"]
meshlet_processor = ["meshlet", "dep:meshopt", "dep:metis", "dep:itertools"]
[dependencies]
# bevy
@ -37,18 +43,25 @@ bevy_utils = { path = "../bevy_utils", version = "0.14.0-dev" }
bevy_window = { path = "../bevy_window", version = "0.14.0-dev" }
bevy_derive = { path = "../bevy_derive", version = "0.14.0-dev" }
# other
meshopt = { version = "0.2", optional = true }
thiserror = { version = "1", optional = true }
bitflags = "2.3"
fixedbitset = "0.5"
# meshlet
lz4_flex = { version = "0.11", default-features = false, features = [
"frame",
], optional = true }
serde = { version = "1", features = ["derive", "rc"], optional = true }
bincode = { version = "1", optional = true }
thiserror = { version = "1", optional = true }
range-alloc = { version = "0.1", optional = true }
meshopt = { version = "0.2", optional = true }
metis = { version = "0.2", optional = true }
itertools = { version = "0.12", optional = true }
# direct dependency required for derive macro
bytemuck = { version = "1", features = ["derive", "must_cast"] }
radsort = "0.1"
smallvec = "1.6"
serde = { version = "1", features = ["derive", "rc"] }
bincode = { version = "1", optional = true }
range-alloc = { version = "0.1", optional = true }
nonmax = "0.5"
static_assertions = "1"

View file

@ -1,7 +1,7 @@
// FIXME(3492): remove once docs are ready
#![allow(missing_docs)]
#![cfg_attr(docsrs, feature(doc_auto_cfg))]
#![forbid(unsafe_code)]
#![deny(unsafe_code)]
#![doc(
html_logo_url = "https://bevyengine.org/assets/icon.png",
html_favicon_url = "https://bevyengine.org/assets/icon.png"

View file

@ -1,13 +1,17 @@
use bevy_asset::{
io::{Reader, Writer},
io::{AsyncReadAndSeek, Reader, Writer},
saver::{AssetSaver, SavedAsset},
Asset, AssetLoader, AsyncReadExt, AsyncWriteExt, LoadContext,
};
use bevy_math::Vec3;
use bevy_reflect::TypePath;
use bytemuck::{Pod, Zeroable};
use lz4_flex::frame::{FrameDecoder, FrameEncoder};
use serde::{Deserialize, Serialize};
use std::sync::Arc;
use std::{io::Cursor, sync::Arc};
/// The current version of the [`MeshletMesh`] asset format.
pub const MESHLET_MESH_ASSET_VERSION: u64 = 0;
/// A mesh that has been pre-processed into multiple small clusters of triangles called meshlets.
///
@ -25,8 +29,8 @@ use std::sync::Arc;
/// See also [`super::MaterialMeshletMeshBundle`] and [`super::MeshletPlugin`].
#[derive(Asset, TypePath, Serialize, Deserialize, Clone)]
pub struct MeshletMesh {
/// The total amount of triangles summed across all meshlets in the mesh.
pub total_meshlet_triangles: u64,
/// The total amount of triangles summed across all LOD 0 meshlets in the mesh.
pub worst_case_meshlet_triangles: u64,
/// Raw vertex data bytes for the overall mesh.
pub vertex_data: Arc<[u8]>,
/// Indices into `vertex_data`.
@ -35,8 +39,8 @@ pub struct MeshletMesh {
pub indices: Arc<[u8]>,
/// The list of meshlets making up this mesh.
pub meshlets: Arc<[Meshlet]>,
/// A list of spherical bounding volumes, 1 per meshlet.
pub meshlet_bounding_spheres: Arc<[MeshletBoundingSphere]>,
/// Spherical bounding volumes.
pub bounding_spheres: Arc<[MeshletBoundingSpheres]>,
}
/// A single meshlet within a [`MeshletMesh`].
@ -51,7 +55,19 @@ pub struct Meshlet {
pub triangle_count: u32,
}
/// A spherical bounding volume used for culling a [`Meshlet`].
/// Bounding spheres used for culling and choosing level of detail for a [`Meshlet`].
#[derive(Serialize, Deserialize, Copy, Clone, Pod, Zeroable)]
#[repr(C)]
pub struct MeshletBoundingSpheres {
/// The bounding sphere used for frustum and occlusion culling for this meshlet.
pub self_culling: MeshletBoundingSphere,
/// The bounding sphere used for determining if this meshlet is at the correct level of detail for a given view.
pub self_lod: MeshletBoundingSphere,
/// The bounding sphere used for determining if this meshlet's parent is at the correct level of detail for a given view.
pub parent_lod: MeshletBoundingSphere,
}
/// A spherical bounding volume used for a [`Meshlet`].
#[derive(Serialize, Deserialize, Copy, Clone, Pod, Zeroable)]
#[repr(C)]
pub struct MeshletBoundingSphere {
@ -65,7 +81,7 @@ pub struct MeshletMeshSaverLoad;
impl AssetLoader for MeshletMeshSaverLoad {
type Asset = MeshletMesh;
type Settings = ();
type Error = bincode::Error;
type Error = MeshletMeshSaveOrLoadError;
async fn load<'a>(
&'a self,
@ -73,9 +89,16 @@ impl AssetLoader for MeshletMeshSaverLoad {
_settings: &'a Self::Settings,
_load_context: &'a mut LoadContext<'_>,
) -> Result<Self::Asset, Self::Error> {
let version = read_u64(reader).await?;
if version != MESHLET_MESH_ASSET_VERSION {
return Err(MeshletMeshSaveOrLoadError::WrongVersion { found: version });
}
let mut bytes = Vec::new();
reader.read_to_end(&mut bytes).await?;
bincode::deserialize(&bytes)
let asset = bincode::deserialize_from(FrameDecoder::new(Cursor::new(bytes)))?;
Ok(asset)
}
fn extensions(&self) -> &[&str] {
@ -87,7 +110,7 @@ impl AssetSaver for MeshletMeshSaverLoad {
type Asset = MeshletMesh;
type Settings = ();
type OutputLoader = Self;
type Error = bincode::Error;
type Error = MeshletMeshSaveOrLoadError;
async fn save<'a>(
&'a self,
@ -95,8 +118,36 @@ impl AssetSaver for MeshletMeshSaverLoad {
asset: SavedAsset<'a, Self::Asset>,
_settings: &'a Self::Settings,
) -> Result<(), Self::Error> {
let bytes = bincode::serialize(asset.get())?;
writer
.write_all(&MESHLET_MESH_ASSET_VERSION.to_le_bytes())
.await?;
let mut bytes = Vec::new();
let mut sync_writer = FrameEncoder::new(&mut bytes);
bincode::serialize_into(&mut sync_writer, asset.get())?;
sync_writer.finish()?;
writer.write_all(&bytes).await?;
Ok(())
}
}
#[derive(thiserror::Error, Debug)]
pub enum MeshletMeshSaveOrLoadError {
#[error("expected asset version {MESHLET_MESH_ASSET_VERSION} but found version {found}")]
WrongVersion { found: u64 },
#[error("failed to serialize or deserialize asset data")]
SerializationOrDeserialization(#[from] bincode::Error),
#[error("failed to compress or decompress asset data")]
CompressionOrDecompression(#[from] lz4_flex::frame::Error),
#[error("failed to read or write asset data")]
Io(#[from] std::io::Error),
}
async fn read_u64(
reader: &mut (dyn AsyncReadAndSeek + Sync + Send + Unpin),
) -> Result<u64, bincode::Error> {
let mut bytes = [0u8; 8];
reader.read_exact(&mut bytes).await?;
Ok(u64::from_le_bytes(bytes))
}

View file

@ -20,19 +20,34 @@
@compute
@workgroup_size(128, 1, 1) // 128 threads per workgroup, 1 instanced meshlet per thread
fn cull_meshlets(@builtin(global_invocation_id) cluster_id: vec3<u32>) {
// Fetch the instanced meshlet data
// Fetch the instance data and check for instance culling
if cluster_id.x >= arrayLength(&meshlet_thread_meshlet_ids) { return; }
let instance_id = meshlet_thread_instance_ids[cluster_id.x];
if should_cull_instance(instance_id) {
return;
}
// Fetch other meshlet data
let meshlet_id = meshlet_thread_meshlet_ids[cluster_id.x];
let bounding_sphere = meshlet_bounding_spheres[meshlet_id];
let instance_uniform = meshlet_instance_uniforms[instance_id];
let model = affine3_to_square(instance_uniform.model);
let model_scale = max(length(model[0]), max(length(model[1]), length(model[2])));
let bounding_sphere_center = model * vec4(bounding_sphere.center, 1.0);
let bounding_sphere_radius = model_scale * bounding_sphere.radius;
let bounding_spheres = meshlet_bounding_spheres[meshlet_id];
// Calculate view-space LOD bounding sphere for the meshlet
let lod_bounding_sphere_center = model * vec4(bounding_spheres.self_lod.center, 1.0);
let lod_bounding_sphere_radius = model_scale * bounding_spheres.self_lod.radius;
let lod_bounding_sphere_center_view_space = (view.inverse_view * vec4(lod_bounding_sphere_center.xyz, 1.0)).xyz;
// Calculate view-space LOD bounding sphere for the meshlet's parent
let parent_lod_bounding_sphere_center = model * vec4(bounding_spheres.parent_lod.center, 1.0);
let parent_lod_bounding_sphere_radius = model_scale * bounding_spheres.parent_lod.radius;
let parent_lod_bounding_sphere_center_view_space = (view.inverse_view * vec4(parent_lod_bounding_sphere_center.xyz, 1.0)).xyz;
// Check LOD cut (meshlet error imperceptible, and parent error not imperceptible)
let lod_is_ok = lod_error_is_imperceptible(lod_bounding_sphere_center_view_space, lod_bounding_sphere_radius);
let parent_lod_is_ok = lod_error_is_imperceptible(parent_lod_bounding_sphere_center_view_space, parent_lod_bounding_sphere_radius);
if !lod_is_ok || parent_lod_is_ok { return; }
// In the first pass, operate only on the clusters visible last frame. In the second pass, operate on all clusters.
#ifdef MESHLET_SECOND_CULLING_PASS
@ -42,18 +57,22 @@ fn cull_meshlets(@builtin(global_invocation_id) cluster_id: vec3<u32>) {
if !meshlet_visible { return; }
#endif
// Calculate world-space culling bounding sphere for the cluster
let culling_bounding_sphere_center = model * vec4(bounding_spheres.self_culling.center, 1.0);
let culling_bounding_sphere_radius = model_scale * bounding_spheres.self_culling.radius;
// Frustum culling
// TODO: Faster method from https://vkguide.dev/docs/gpudriven/compute_culling/#frustum-culling-function
for (var i = 0u; i < 6u; i++) {
if !meshlet_visible { break; }
meshlet_visible &= dot(view.frustum[i], bounding_sphere_center) > -bounding_sphere_radius;
meshlet_visible &= dot(view.frustum[i], culling_bounding_sphere_center) > -culling_bounding_sphere_radius;
}
#ifdef MESHLET_SECOND_CULLING_PASS
// In the second culling pass, cull against the depth pyramid generated from the first pass
if meshlet_visible {
let bounding_sphere_center_view_space = (view.inverse_view * vec4(bounding_sphere_center.xyz, 1.0)).xyz;
let aabb = project_view_space_sphere_to_screen_space_aabb(bounding_sphere_center_view_space, bounding_sphere_radius);
let culling_bounding_sphere_center_view_space = (view.inverse_view * vec4(culling_bounding_sphere_center.xyz, 1.0)).xyz;
let aabb = project_view_space_sphere_to_screen_space_aabb(culling_bounding_sphere_center_view_space, culling_bounding_sphere_radius);
// Halve the AABB size because the first depth mip resampling pass cut the full screen resolution into a power of two conservatively
let depth_pyramid_size_mip_0 = vec2<f32>(textureDimensions(depth_pyramid, 0)) * 0.5;
@ -71,11 +90,11 @@ fn cull_meshlets(@builtin(global_invocation_id) cluster_id: vec3<u32>) {
let occluder_depth = min(min(depth_quad_a, depth_quad_b), min(depth_quad_c, depth_quad_d));
if view.projection[3][3] == 1.0 {
// Orthographic
let sphere_depth = view.projection[3][2] + (bounding_sphere_center_view_space.z + bounding_sphere_radius) * view.projection[2][2];
let sphere_depth = view.projection[3][2] + (culling_bounding_sphere_center_view_space.z + culling_bounding_sphere_radius) * view.projection[2][2];
meshlet_visible &= sphere_depth >= occluder_depth;
} else {
// Perspective
let sphere_depth = -view.projection[3][2] / (bounding_sphere_center_view_space.z + bounding_sphere_radius);
let sphere_depth = -view.projection[3][2] / (culling_bounding_sphere_center_view_space.z + culling_bounding_sphere_radius);
meshlet_visible &= sphere_depth >= occluder_depth;
}
}
@ -86,6 +105,16 @@ fn cull_meshlets(@builtin(global_invocation_id) cluster_id: vec3<u32>) {
atomicOr(&meshlet_occlusion[cluster_id.x / 32u], occlusion_bit);
}
// https://stackoverflow.com/questions/21648630/radius-of-projected-sphere-in-screen-space/21649403#21649403
fn lod_error_is_imperceptible(cp: vec3<f32>, r: f32) -> bool {
let d2 = dot(cp, cp);
let r2 = r * r;
let sphere_diameter_uv = view.projection[0][0] * r / sqrt(d2 - r2);
let view_size = f32(max(view.viewport.z, view.viewport.w));
let sphere_diameter_pixels = sphere_diameter_uv * view_size;
return sphere_diameter_pixels < 1.0;
}
// https://zeux.io/2023/01/12/approximate-projected-bounds
fn project_view_space_sphere_to_screen_space_aabb(cp: vec3<f32>, r: f32) -> vec4<f32> {
let inv_width = view.projection[0][0] * 0.5;

View file

@ -1,10 +1,17 @@
use super::asset::{Meshlet, MeshletBoundingSphere, MeshletMesh};
use super::asset::{Meshlet, MeshletBoundingSphere, MeshletBoundingSpheres, MeshletMesh};
use bevy_render::{
mesh::{Indices, Mesh},
render_resource::PrimitiveTopology,
};
use meshopt::{build_meshlets, compute_meshlet_bounds_decoder, VertexDataAdapter};
use std::borrow::Cow;
use bevy_utils::{HashMap, HashSet};
use itertools::Itertools;
use meshopt::{
build_meshlets, compute_cluster_bounds, compute_meshlet_bounds,
ffi::{meshopt_Bounds, meshopt_optimizeMeshlet},
simplify, simplify_scale, Meshlets, SimplifyOptions, VertexDataAdapter,
};
use metis::Graph;
use std::{borrow::Cow, ops::Range};
impl MeshletMesh {
/// Process a [`Mesh`] to generate a [`MeshletMesh`].
@ -19,80 +26,342 @@ impl MeshletMesh {
/// 3. Have the exact following set of vertex attributes: `{POSITION, NORMAL, UV_0, TANGENT}`
pub fn from_mesh(mesh: &Mesh) -> Result<Self, MeshToMeshletMeshConversionError> {
// Validate mesh format
if mesh.primitive_topology() != PrimitiveTopology::TriangleList {
return Err(MeshToMeshletMeshConversionError::WrongMeshPrimitiveTopology);
}
if mesh.attributes().map(|(id, _)| id).ne([
Mesh::ATTRIBUTE_POSITION.id,
Mesh::ATTRIBUTE_NORMAL.id,
Mesh::ATTRIBUTE_UV_0.id,
Mesh::ATTRIBUTE_TANGENT.id,
]) {
return Err(MeshToMeshletMeshConversionError::WrongMeshVertexAttributes);
}
let indices = match mesh.indices() {
Some(Indices::U32(indices)) => Cow::Borrowed(indices.as_slice()),
Some(Indices::U16(indices)) => indices.iter().map(|i| *i as u32).collect(),
_ => return Err(MeshToMeshletMeshConversionError::MeshMissingIndices),
};
let indices = validate_input_mesh(mesh)?;
// Split the mesh into an initial list of meshlets (LOD 0)
let vertex_buffer = mesh.get_vertex_buffer_data();
let vertices =
VertexDataAdapter::new(&vertex_buffer, mesh.get_vertex_size() as usize, 0).unwrap();
// Split the mesh into meshlets
let meshopt_meshlets = build_meshlets(&indices, &vertices, 64, 64, 0.0);
// Calculate meshlet bounding spheres
let meshlet_bounding_spheres = meshopt_meshlets
let vertex_stride = mesh.get_vertex_size() as usize;
let vertices = VertexDataAdapter::new(&vertex_buffer, vertex_stride, 0).unwrap();
let mut meshlets = compute_meshlets(&indices, &vertices);
let mut bounding_spheres = meshlets
.iter()
.map(|meshlet| {
compute_meshlet_bounds_decoder(
meshlet,
mesh.attribute(Mesh::ATTRIBUTE_POSITION)
.unwrap()
.as_float3()
.unwrap(),
)
.map(|meshlet| compute_meshlet_bounds(meshlet, &vertices))
.map(convert_meshlet_bounds)
.map(|bounding_sphere| MeshletBoundingSpheres {
self_culling: bounding_sphere,
self_lod: MeshletBoundingSphere {
center: bounding_sphere.center,
radius: 0.0,
},
parent_lod: MeshletBoundingSphere {
center: bounding_sphere.center,
radius: f32::MAX,
},
})
.map(|bounds| MeshletBoundingSphere {
center: bounds.center.into(),
radius: bounds.radius,
})
.collect();
.collect::<Vec<_>>();
let worst_case_meshlet_triangles = meshlets
.meshlets
.iter()
.map(|m| m.triangle_count as u64)
.sum();
// Assemble into the final asset
let mut total_meshlet_triangles = 0;
let meshlets = meshopt_meshlets
// Build further LODs
let mut simplification_queue = 0..meshlets.len();
let mut lod_level = 1;
while simplification_queue.len() > 1 {
// For each meshlet build a set of triangle edges
let triangle_edges_per_meshlet =
collect_triangle_edges_per_meshlet(simplification_queue.clone(), &meshlets);
// For each meshlet build a list of connected meshlets (meshlets that share a triangle edge)
let connected_meshlets_per_meshlet =
find_connected_meshlets(simplification_queue.clone(), &triangle_edges_per_meshlet);
// Group meshlets into roughly groups of 4, grouping meshlets with a high number of shared edges
// http://glaros.dtc.umn.edu/gkhome/fetch/sw/metis/manual.pdf
let groups = group_meshlets(
simplification_queue.clone(),
&connected_meshlets_per_meshlet,
);
let next_lod_start = meshlets.len();
for group_meshlets in groups.values().filter(|group| group.len() > 1) {
// Simplify the group to ~50% triangle count
let Some((simplified_group_indices, mut group_error)) =
simplify_meshlet_groups(group_meshlets, &meshlets, &vertices, lod_level)
else {
continue;
};
// Add the maximum child error to the parent error to make parent error cumulative from LOD 0
// (we're currently building the parent from its children)
group_error += group_meshlets.iter().fold(group_error, |acc, meshlet_id| {
acc.max(bounding_spheres[*meshlet_id].self_lod.radius)
});
// Build a new LOD bounding sphere for the simplified group as a whole
let mut group_bounding_sphere = convert_meshlet_bounds(compute_cluster_bounds(
&simplified_group_indices,
&vertices,
));
group_bounding_sphere.radius = group_error;
// For each meshlet in the group set their parent LOD bounding sphere to that of the simplified group
for meshlet_id in group_meshlets {
bounding_spheres[*meshlet_id].parent_lod = group_bounding_sphere;
}
// Build new meshlets using the simplified group
let new_meshlets_count = split_simplified_groups_into_new_meshlets(
&simplified_group_indices,
&vertices,
&mut meshlets,
);
// Calculate the culling bounding sphere for the new meshlets and set their LOD bounding spheres
let new_meshlet_ids = (meshlets.len() - new_meshlets_count)..meshlets.len();
bounding_spheres.extend(
new_meshlet_ids
.map(|meshlet_id| {
compute_meshlet_bounds(meshlets.get(meshlet_id), &vertices)
})
.map(convert_meshlet_bounds)
.map(|bounding_sphere| MeshletBoundingSpheres {
self_culling: bounding_sphere,
self_lod: group_bounding_sphere,
parent_lod: MeshletBoundingSphere {
center: group_bounding_sphere.center,
radius: f32::MAX,
},
}),
);
}
simplification_queue = next_lod_start..meshlets.len();
lod_level += 1;
}
// Convert meshopt_Meshlet data to a custom format
let bevy_meshlets = meshlets
.meshlets
.into_iter()
.map(|m| {
total_meshlet_triangles += m.triangle_count as u64;
Meshlet {
start_vertex_id: m.vertex_offset,
start_index_id: m.triangle_offset,
triangle_count: m.triangle_count,
}
.map(|m| Meshlet {
start_vertex_id: m.vertex_offset,
start_index_id: m.triangle_offset,
triangle_count: m.triangle_count,
})
.collect();
Ok(Self {
total_meshlet_triangles,
worst_case_meshlet_triangles,
vertex_data: vertex_buffer.into(),
vertex_ids: meshopt_meshlets.vertices.into(),
indices: meshopt_meshlets.triangles.into(),
meshlets,
meshlet_bounding_spheres,
vertex_ids: meshlets.vertices.into(),
indices: meshlets.triangles.into(),
meshlets: bevy_meshlets,
bounding_spheres: bounding_spheres.into(),
})
}
}
fn validate_input_mesh(mesh: &Mesh) -> Result<Cow<'_, [u32]>, MeshToMeshletMeshConversionError> {
if mesh.primitive_topology() != PrimitiveTopology::TriangleList {
return Err(MeshToMeshletMeshConversionError::WrongMeshPrimitiveTopology);
}
if mesh.attributes().map(|(id, _)| id).ne([
Mesh::ATTRIBUTE_POSITION.id,
Mesh::ATTRIBUTE_NORMAL.id,
Mesh::ATTRIBUTE_UV_0.id,
Mesh::ATTRIBUTE_TANGENT.id,
]) {
return Err(MeshToMeshletMeshConversionError::WrongMeshVertexAttributes);
}
match mesh.indices() {
Some(Indices::U32(indices)) => Ok(Cow::Borrowed(indices.as_slice())),
Some(Indices::U16(indices)) => Ok(indices.iter().map(|i| *i as u32).collect()),
_ => Err(MeshToMeshletMeshConversionError::MeshMissingIndices),
}
}
fn compute_meshlets(indices: &[u32], vertices: &VertexDataAdapter) -> Meshlets {
let mut meshlets = build_meshlets(indices, vertices, 64, 64, 0.0);
for meshlet in &mut meshlets.meshlets {
#[allow(unsafe_code)]
#[allow(clippy::undocumented_unsafe_blocks)]
unsafe {
meshopt_optimizeMeshlet(
&mut meshlets.vertices[meshlet.vertex_offset as usize],
&mut meshlets.triangles[meshlet.triangle_offset as usize],
meshlet.triangle_count as usize,
meshlet.vertex_count as usize,
);
}
}
meshlets
}
fn collect_triangle_edges_per_meshlet(
simplification_queue: Range<usize>,
meshlets: &Meshlets,
) -> HashMap<usize, HashSet<(u32, u32)>> {
let mut triangle_edges_per_meshlet = HashMap::new();
for meshlet_id in simplification_queue {
let meshlet = meshlets.get(meshlet_id);
let meshlet_triangle_edges = triangle_edges_per_meshlet
.entry(meshlet_id)
.or_insert(HashSet::new());
for i in meshlet.triangles.chunks(3) {
let v0 = meshlet.vertices[i[0] as usize];
let v1 = meshlet.vertices[i[1] as usize];
let v2 = meshlet.vertices[i[2] as usize];
meshlet_triangle_edges.insert((v0.min(v1), v0.max(v1)));
meshlet_triangle_edges.insert((v0.min(v2), v0.max(v2)));
meshlet_triangle_edges.insert((v1.min(v2), v1.max(v2)));
}
}
triangle_edges_per_meshlet
}
fn find_connected_meshlets(
simplification_queue: Range<usize>,
triangle_edges_per_meshlet: &HashMap<usize, HashSet<(u32, u32)>>,
) -> HashMap<usize, Vec<(usize, usize)>> {
let mut connected_meshlets_per_meshlet = HashMap::new();
for meshlet_id in simplification_queue.clone() {
connected_meshlets_per_meshlet.insert(meshlet_id, Vec::new());
}
for (meshlet_id1, meshlet_id2) in simplification_queue.tuple_combinations() {
let shared_edge_count = triangle_edges_per_meshlet[&meshlet_id1]
.intersection(&triangle_edges_per_meshlet[&meshlet_id2])
.count();
if shared_edge_count != 0 {
connected_meshlets_per_meshlet
.get_mut(&meshlet_id1)
.unwrap()
.push((meshlet_id2, shared_edge_count));
connected_meshlets_per_meshlet
.get_mut(&meshlet_id2)
.unwrap()
.push((meshlet_id1, shared_edge_count));
}
}
connected_meshlets_per_meshlet
}
fn group_meshlets(
simplification_queue: Range<usize>,
connected_meshlets_per_meshlet: &HashMap<usize, Vec<(usize, usize)>>,
) -> HashMap<i32, Vec<usize>> {
let mut xadj = Vec::with_capacity(simplification_queue.len() + 1);
let mut adjncy = Vec::new();
let mut adjwgt = Vec::new();
for meshlet_id in simplification_queue.clone() {
xadj.push(adjncy.len() as i32);
for (connected_meshlet_id, shared_edge_count) in
connected_meshlets_per_meshlet[&meshlet_id].iter().copied()
{
adjncy.push((connected_meshlet_id - simplification_queue.start) as i32);
adjwgt.push(shared_edge_count as i32);
}
}
xadj.push(adjncy.len() as i32);
let mut group_per_meshlet = vec![0; simplification_queue.len()];
let partition_count = (simplification_queue.len().div_ceil(4)) as i32;
Graph::new(1, partition_count, &xadj, &adjncy)
.unwrap()
.set_adjwgt(&adjwgt)
.part_kway(&mut group_per_meshlet)
.unwrap();
let mut groups = HashMap::new();
for (i, meshlet_group) in group_per_meshlet.into_iter().enumerate() {
groups
.entry(meshlet_group)
.or_insert(Vec::new())
.push(i + simplification_queue.start);
}
groups
}
fn simplify_meshlet_groups(
group_meshlets: &[usize],
meshlets: &Meshlets,
vertices: &VertexDataAdapter<'_>,
lod_level: u32,
) -> Option<(Vec<u32>, f32)> {
// Build a new index buffer into the mesh vertex data by combining all meshlet data in the group
let mut group_indices = Vec::new();
for meshlet_id in group_meshlets {
let meshlet = meshlets.get(*meshlet_id);
for meshlet_index in meshlet.triangles {
group_indices.push(meshlet.vertices[*meshlet_index as usize]);
}
}
// Allow more deformation for high LOD levels (1% at LOD 1, 10% at LOD 20+)
let t = (lod_level - 1) as f32 / 19.0;
let target_error = 0.1 * t + 0.01 * (1.0 - t);
// Simplify the group to ~50% triangle count
// TODO: Use simplify_with_locks()
let mut error = 0.0;
let simplified_group_indices = simplify(
&group_indices,
vertices,
group_indices.len() / 2,
target_error,
SimplifyOptions::LockBorder,
Some(&mut error),
);
// Check if we were able to simplify to at least 65% triangle count
if simplified_group_indices.len() as f32 / group_indices.len() as f32 > 0.65 {
return None;
}
// Convert error to object-space and convert from diameter to radius
error *= simplify_scale(vertices) * 0.5;
Some((simplified_group_indices, error))
}
fn split_simplified_groups_into_new_meshlets(
simplified_group_indices: &[u32],
vertices: &VertexDataAdapter<'_>,
meshlets: &mut Meshlets,
) -> usize {
let simplified_meshlets = compute_meshlets(simplified_group_indices, vertices);
let new_meshlets_count = simplified_meshlets.len();
let vertex_offset = meshlets.vertices.len() as u32;
let triangle_offset = meshlets.triangles.len() as u32;
meshlets
.vertices
.extend_from_slice(&simplified_meshlets.vertices);
meshlets
.triangles
.extend_from_slice(&simplified_meshlets.triangles);
meshlets
.meshlets
.extend(simplified_meshlets.meshlets.into_iter().map(|mut meshlet| {
meshlet.vertex_offset += vertex_offset;
meshlet.triangle_offset += triangle_offset;
meshlet
}));
new_meshlets_count
}
fn convert_meshlet_bounds(bounds: meshopt_Bounds) -> MeshletBoundingSphere {
MeshletBoundingSphere {
center: bounds.center.into(),
radius: bounds.radius,
}
}
/// An error produced by [`MeshletMesh::from_mesh`].
#[derive(thiserror::Error, Debug)]
pub enum MeshToMeshletMeshConversionError {
#[error("Mesh primitive topology was not TriangleList")]
#[error("Mesh primitive topology is not TriangleList")]
WrongMeshPrimitiveTopology,
#[error("Mesh attributes were not {{POSITION, NORMAL, UV_0, TANGENT}}")]
#[error("Mesh attributes are not {{POSITION, NORMAL, UV_0, TANGENT}}")]
WrongMeshVertexAttributes,
#[error("Mesh had no indices")]
#[error("Mesh has no indices")]
MeshMissingIndices,
}

View file

@ -1,4 +1,7 @@
use super::{persistent_buffer::PersistentGpuBuffer, Meshlet, MeshletBoundingSphere, MeshletMesh};
use super::{
asset::{Meshlet, MeshletBoundingSpheres, MeshletMesh},
persistent_buffer::PersistentGpuBuffer,
};
use crate::{
Material, MeshFlags, MeshTransforms, MeshUniform, NotShadowCaster, NotShadowReceiver,
PreviousGlobalTransform, RenderMaterialInstances, ShadowView,
@ -352,7 +355,7 @@ pub fn prepare_meshlet_per_frame_resources(
});
let depth_size = Extent3d {
// If not a power of 2, round down to the nearest power of 2 to ensure depth is conservative
// Round down to the nearest power of 2 to ensure depth is conservative
width: previous_power_of_2(view.viewport.z),
height: previous_power_of_2(view.viewport.w),
depth_or_array_layers: 1,
@ -611,7 +614,7 @@ pub struct MeshletGpuScene {
vertex_ids: PersistentGpuBuffer<Arc<[u32]>>,
indices: PersistentGpuBuffer<Arc<[u8]>>,
meshlets: PersistentGpuBuffer<Arc<[Meshlet]>>,
meshlet_bounding_spheres: PersistentGpuBuffer<Arc<[MeshletBoundingSphere]>>,
meshlet_bounding_spheres: PersistentGpuBuffer<Arc<[MeshletBoundingSpheres]>>,
meshlet_mesh_slices: HashMap<AssetId<MeshletMesh>, ([Range<BufferAddress>; 5], u64)>,
scene_meshlet_count: u32,
@ -840,7 +843,7 @@ impl MeshletGpuScene {
);
let meshlet_bounding_spheres_slice = self
.meshlet_bounding_spheres
.queue_write(Arc::clone(&meshlet_mesh.meshlet_bounding_spheres), ());
.queue_write(Arc::clone(&meshlet_mesh.bounding_spheres), ());
(
[
@ -850,7 +853,7 @@ impl MeshletGpuScene {
meshlets_slice,
meshlet_bounding_spheres_slice,
],
meshlet_mesh.total_meshlet_triangles,
meshlet_mesh.worst_case_meshlet_triangles,
)
};

View file

@ -135,8 +135,6 @@ pub fn prepare_material_meshlet_meshes_main_opaque_pass<M: Material>(
view_key |= MeshPipelineKey::SCREEN_SPACE_AMBIENT_OCCLUSION;
}
// TODO: Lightmaps
view_key |= MeshPipelineKey::from_primitive_topology(PrimitiveTopology::TriangleList);
for material_id in render_material_instances.values() {

View file

@ -32,6 +32,12 @@ struct Meshlet {
triangle_count: u32,
}
struct MeshletBoundingSpheres {
self_culling: MeshletBoundingSphere,
self_lod: MeshletBoundingSphere,
parent_lod: MeshletBoundingSphere,
}
struct MeshletBoundingSphere {
center: vec3<f32>,
radius: f32,
@ -46,7 +52,7 @@ struct DrawIndirectArgs {
#ifdef MESHLET_CULLING_PASS
@group(0) @binding(0) var<storage, read> meshlet_thread_meshlet_ids: array<u32>; // Per cluster (instance of a meshlet)
@group(0) @binding(1) var<storage, read> meshlet_bounding_spheres: array<MeshletBoundingSphere>; // Per asset meshlet
@group(0) @binding(1) var<storage, read> meshlet_bounding_spheres: array<MeshletBoundingSpheres>; // Per asset meshlet
@group(0) @binding(2) var<storage, read> meshlet_thread_instance_ids: array<u32>; // Per cluster (instance of a meshlet)
@group(0) @binding(3) var<storage, read> meshlet_instance_uniforms: array<Mesh>; // Per entity instance
@group(0) @binding(4) var<storage, read> meshlet_view_instance_visibility: array<u32>; // 1 bit per entity instance, packed as a bitmask

View file

@ -30,12 +30,11 @@ pub(crate) use self::{
},
};
pub use self::asset::{Meshlet, MeshletBoundingSphere, MeshletMesh};
pub use self::asset::*;
#[cfg(feature = "meshlet_processor")]
pub use self::from_mesh::MeshToMeshletMeshConversionError;
use self::{
asset::MeshletMeshSaverLoad,
gpu_scene::{
extract_meshlet_meshes, perform_pending_meshlet_mesh_writes,
prepare_meshlet_per_frame_resources, prepare_meshlet_view_bind_groups,
@ -99,6 +98,7 @@ const MESHLET_MESH_MATERIAL_SHADER_HANDLE: Handle<Shader> =
/// Additionally, occlusion culling can eliminate meshlets that would cause overdraw.
/// * Much more efficient batching. All geometry can be rasterized in a single indirect draw.
/// * Scales better with large amounts of dense geometry and overdraw. Bevy's standard renderer will bottleneck sooner.
/// * Near-seamless level of detail (LOD).
/// * Much greater base overhead. Rendering will be slower than Bevy's standard renderer with small amounts of geometry and overdraw.
/// * Much greater memory usage.
/// * Requires preprocessing meshes. See [`MeshletMesh`] for details.
@ -196,17 +196,18 @@ impl Plugin for MeshletPlugin {
.add_render_graph_edges(
Core3d,
(
NodeMeshlet::VisibilityBufferRasterPass,
// TODO: Meshlet VisibilityBufferRaster should be after main pass when not using depth prepass
NodePbr::ShadowPass,
NodeMeshlet::Prepass,
NodeMeshlet::DeferredPrepass,
Node3d::Prepass,
Node3d::DeferredPrepass,
NodeMeshlet::VisibilityBufferRasterPass,
NodeMeshlet::Prepass,
NodeMeshlet::DeferredPrepass,
Node3d::CopyDeferredLightingId,
Node3d::EndPrepasses,
Node3d::StartMainPass,
NodeMeshlet::MainOpaquePass,
Node3d::MainOpaquePass,
NodeMeshlet::MainOpaquePass,
Node3d::EndMainPass,
),
)
@ -282,6 +283,7 @@ fn configure_meshlet_views(
.entity(entity)
.insert(MeshletViewMaterialsMainOpaquePass::default());
} else {
// TODO: Should we add both Prepass and DeferredGBufferPrepass materials here, and in other systems/nodes?
commands.entity(entity).insert((
MeshletViewMaterialsMainOpaquePass::default(),
MeshletViewMaterialsPrepass::default(),

View file

@ -1,4 +1,7 @@
use super::{persistent_buffer::PersistentGpuBufferable, Meshlet, MeshletBoundingSphere};
use super::{
asset::{Meshlet, MeshletBoundingSpheres},
persistent_buffer::PersistentGpuBufferable,
};
use std::{mem::size_of, sync::Arc};
const MESHLET_VERTEX_SIZE_IN_BYTES: u32 = 48;
@ -62,11 +65,11 @@ impl PersistentGpuBufferable for Arc<[Meshlet]> {
}
}
impl PersistentGpuBufferable for Arc<[MeshletBoundingSphere]> {
impl PersistentGpuBufferable for Arc<[MeshletBoundingSpheres]> {
type Metadata = ();
fn size_in_bytes(&self) -> usize {
self.len() * size_of::<MeshletBoundingSphere>()
self.len() * size_of::<MeshletBoundingSpheres>()
}
fn write_bytes_le(&self, _: Self::Metadata, buffer_slice: &mut [u8]) {

View file

@ -1,22 +1,29 @@
//! Meshlet rendering for dense high-poly scenes (experimental).
// Note: This example showcases the meshlet API, but is not the type of scene that would benefit from using meshlets.
#[path = "../helpers/camera_controller.rs"]
mod camera_controller;
use bevy::{
pbr::{
experimental::meshlet::{MaterialMeshletMeshBundle, MeshletMesh, MeshletPlugin},
experimental::meshlet::{MaterialMeshletMeshBundle, MeshletPlugin},
CascadeShadowConfigBuilder, DirectionalLightShadowMap,
},
prelude::*,
render::render_resource::AsBindGroup,
};
use camera_controller::{CameraController, CameraControllerPlugin};
use std::f32::consts::PI;
use std::{f32::consts::PI, path::Path, process::ExitCode};
// Note: This example showcases the meshlet API, but is not the type of scene that would benefit from using meshlets.
const ASSET_URL: &str = "https://github.com/JMS55/bevy_meshlet_asset/blob/bd869887bc5c9c6e74e353f657d342bef84bacd8/bunny.meshlet_mesh";
fn main() -> ExitCode {
if !Path::new("./assets/models/bunny.meshlet_mesh").exists() {
println!("ERROR: Asset at path <bevy>/assets/models/bunny.meshlet_mesh is missing. Please download it from {ASSET_URL}");
return ExitCode::FAILURE;
}
fn main() {
App::new()
.insert_resource(DirectionalLightShadowMap { size: 4096 })
.add_plugins((
@ -26,8 +33,9 @@ fn main() {
CameraControllerPlugin,
))
.add_systems(Startup, setup)
.add_systems(Update, draw_bounding_spheres)
.run();
ExitCode::SUCCESS
}
fn setup(
@ -37,8 +45,6 @@ fn setup(
mut debug_materials: ResMut<Assets<MeshletDebugMaterial>>,
mut meshes: ResMut<Assets<Mesh>>,
) {
info!("\nMeshlet Controls:\n Space - Toggle bounding spheres");
commands.spawn((
Camera3dBundle {
transform: Transform::from_translation(Vec3::new(1.8, 0.4, -0.1))
@ -61,7 +67,7 @@ fn setup(
},
cascade_shadow_config: CascadeShadowConfigBuilder {
num_cascades: 1,
maximum_distance: 5.0,
maximum_distance: 15.0,
..default()
}
.build(),
@ -125,54 +131,6 @@ fn setup(
});
}
#[allow(clippy::too_many_arguments)]
fn draw_bounding_spheres(
query: Query<(&Handle<MeshletMesh>, &Transform), With<Handle<MeshletDebugMaterial>>>,
debug: Query<&MeshletBoundingSpheresDebug>,
camera: Query<&Transform, With<Camera>>,
mut commands: Commands,
meshlets: Res<Assets<MeshletMesh>>,
mut gizmos: Gizmos,
keys: Res<ButtonInput<KeyCode>>,
mut should_draw: Local<bool>,
) {
if keys.just_pressed(KeyCode::Space) {
*should_draw = !*should_draw;
}
match debug.get_single() {
Ok(meshlet_debug) if *should_draw => {
let camera_pos = camera.single().translation;
for circle in &meshlet_debug.circles {
gizmos.circle(
circle.0,
Dir3::new(camera_pos - circle.0).unwrap(),
circle.1,
Color::BLACK,
);
}
}
Err(_) => {
if let Some((handle, transform)) = query.iter().last() {
if let Some(meshlets) = meshlets.get(handle) {
let mut circles = Vec::new();
for bounding_sphere in meshlets.meshlet_bounding_spheres.iter() {
let center = transform.transform_point(bounding_sphere.center);
circles.push((center, transform.scale.x * bounding_sphere.radius));
}
commands.spawn(MeshletBoundingSpheresDebug { circles });
}
}
}
_ => {}
}
}
#[derive(Component)]
struct MeshletBoundingSpheresDebug {
circles: Vec<(Vec3, f32)>,
}
#[derive(Asset, TypePath, AsBindGroup, Clone, Default)]
struct MeshletDebugMaterial {
_dummy: (),