Meshlet BVH Culling #19318

atlv24 · 2025-05-21T09:34:16Z

Objective

Merge @SparkyPotato 's efforts to implement BVH-accelerated meshlet culling.

Solution

Add hot reloading support
Fix near-plane overculling
Fix hzb sampling
Fix orthographic error metric

Testing

Meshlet example, Nsight, hot-reloading and careful thinking

Readd software raster

crates/bevy_pbr/src/meshlet/meshlet_cull_shared.wgsl

SparkyPotato · 2025-05-26T20:11:34Z

crates/bevy_pbr/src/meshlet/meshlet_cull_shared.wgsl

+        // let world_sphere_radius = lod_sphere.w * world_scale;
+        // let norm_error = simplification_error / world_sphere_radius * 0.25;
+        // return norm_error * view.viewport.w < 1.0;
+        let world_error = simplification_error * world_scale;
+        let proj = projection[1][1];
+        let height = 2.0 / proj;
+        let norm_error = world_error / height;
+        return norm_error * view.viewport.w < 1.0;


Unsure how correct my change (from the commented code) is, though it seems to work.

JMS55

Ok did my first pass of this.

A lot of it I just skimmed because there's way too much code for me to verify line by line, so I'm just going to trust a lot of the BVH traversal stuff and resource handling is correct.

Here's what I still need to verify:

from_mesh() code
Occlusion culling changes
LOD selection changes
Try it out and test perf (my current computer doesn't support 64bit atomics unfortunately)

JMS55 · 2025-06-11T17:19:38Z

crates/bevy_pbr/src/meshlet/asset.rs

-        write_slice(&asset.meshlet_simplification_errors, &mut writer)?;
+        write_slice(&asset.meshlet_cull_data, &mut writer)?;
+        writer.write_all(bytemuck::bytes_of(&asset.aabb))?;
+        writer.write_all(bytemuck::bytes_of(&asset.bvh_depth))?;


Can you move the instance abbb + bvh_depth to before the FrameEncoder?

JMS55 · 2025-06-11T17:20:46Z

crates/bevy_pbr/src/meshlet/mod.rs

-            "remap_1d_to_2d_dispatch.wgsl",
-            Shader::from_wgsl
-        );
+        embedded_asset!(app, "clear_visibility_buffer.wgsl");


Nit: move clear_visibility_buffer to above cull_instances.

JMS55 · 2025-06-11T17:22:07Z

crates/bevy_pbr/src/meshlet/persistent_buffer_impls.rs

+        let base_bvh_node_index = (buffer_offset / size as u64) as u32;
+        for (i, &node) in self.iter().enumerate() {
+            let bytes = bytemuck::cast::<_, [u8; size_of::<BvhNode>()]>(BvhNode {
+                aabbs: core::array::from_fn(|i| {


Can we pull some of this code out into a new function? It's hard to understand imo.

JMS55 · 2025-06-11T17:24:24Z

crates/bevy_pbr/src/meshlet/instance_manager.rs

@@ -233,6 +233,7 @@ pub fn extract_meshlet_mesh_entities(
    }

    // Iterate over every instance
+    // TODO: Switch to change events to not upload every instance every frame.


Yeah better instance handling on the CPU is a big todo. I know @tychedelia is also gonna look into caching material specialization for meshlet instances as well.

JMS55 · 2025-06-11T17:29:04Z

crates/bevy_pbr/src/meshlet/meshlet_cull_shared.wgsl

+#import bevy_render::maths::affine3_to_square
+
+// https://github.com/zeux/meshoptimizer/blob/1e48e96c7e8059321de492865165e9ef071bffba/demo/nanite.cpp#L115
+fn lod_error_is_imperceptible(lod_sphere: vec4<f32>, simplification_error: f32, instance_id: u32) -> bool {


TODO: I need to review this carefully

JMS55 · 2025-06-11T17:35:50Z

crates/bevy_pbr/src/meshlet/resource_manager.rs

-        }
-    };
-    match &mut resource_manager.cluster_meshlet_ids {
+    instance_manager


Did you test that this is equally as fast as upload_storage_buffer?

If I remember correctly, upload_storage_buffer did the same thing internally so I just got rid of it.

I think I might've added that after writing upload_storage_buffer 😅

JMS55 · 2025-06-11T17:40:01Z

crates/bevy_pbr/src/meshlet/meshlet_bindings.wgsl

@@ -48,63 +74,128 @@ struct DrawIndirectArgs {
    first_instance: u32,
 }

+struct InstancedOffset {


Can you add some comments here describing what InstancedOffset is, and show some example code of how you map it to cluster/meshlet/triangle/etc?

JMS55 · 2025-06-11T17:41:16Z

crates/bevy_pbr/src/meshlet/meshlet_cull_shared.wgsl

+    if any((max_texel >> vec2(mip)) > (min_texel >> vec2(mip)) + 3) {
+        mip += 1u;
+    }


Is this the code we can remove when we fix non-pot HiZ?

No, that would be the + 1u (and the clamping) above.

JMS55 · 2025-06-11T17:43:12Z

crates/bevy_pbr/src/meshlet/cull_bvh.wgsl

+}
+
+@compute
+@workgroup_size(128, 1, 1) // 8 threads per node


1 node is handled by 8 threads, so 16 nodes per WG?

JMS55 · 2025-06-11T17:45:04Z

crates/bevy_pbr/src/meshlet/from_mesh.rs

@@ -19,11 +23,13 @@ use meshopt::{
 use metis::{option::Opt, Graph};
 use smallvec::SmallVec;
 use thiserror::Error;
+use tracing::debug_span;



TODO: I need to review this whole file, haven't looked at it yet

JMS55 · 2025-06-11T17:56:01Z

Fix near-plane overculling
Fix hzb sampling
Fix orthographic error metric

What was wrong with each of these?

SparkyPotato and others added 20 commits May 21, 2025 05:28

build bvh

83483ce

uploaded meshes and instances

5911c2c

most scaffolding

66fe2d5

fix shaders somewhat

a8a84cb

something on the screen

df04bc5

some more bugfixes

135202f

Fix never clearing instance_aabbs, formatting

0d87d40

fix flickering and bad bvh cull traversal

6a59614

fix frustum culling

5fd808c

try to render only lod 0

9113ef6

fix meshlet cull not considering all meshlets

b7bba4b

test monotonicity in builder

40b4a65

merge 2 spheres at a time

330a215

fix build

da00224

add occlusion culling

568c5e8

fix occlusion culling

3ed3ea0

feat(meshlet): hot reloading

cbcf716

fix(meshlet): use correct near plane calculation

c867182

fix(meshlet): fix occlusion culling and orthographic error bias

275d655

clippy

5370e14

JMS55 self-requested a review May 21, 2025 14:53

clippy

0e143c6

atlv24 added C-Feature A new feature, making something new possible A-Rendering Drawing game state to the screen S-Needs-Review Needs reviewer attention (from anyone!) to move forward D-Shaders This code uses GPU shader languages labels May 22, 2025

github-project-automation bot added this to Rendering May 22, 2025

atlv24 mentioned this pull request May 26, 2025

Remove Shader weak_handles from bevy_pbr meshlets (with one exception). #19368

Open

SparkyPotato added 2 commits May 26, 2025 17:21

readd sw raster

16e83c7

better orthographic error metric

6d1c71f

Merge pull request #1 from SparkyPotato/bvh-cull-fixed

148a319

Readd software raster

SparkyPotato suggested changes May 26, 2025

View reviewed changes

JMS55 added this to the 0.17 milestone May 27, 2025

feedback

829ba6c

JMS55 suggested changes Jun 11, 2025

View reviewed changes

Uh oh!

Meshlet BVH Culling #19318

Are you sure you want to change the base?

Meshlet BVH Culling #19318

Conversation

atlv24 commented May 21, 2025

Objective

Solution

Testing

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

JMS55 left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

JMS55 commented Jun 11, 2025

Uh oh!

Uh oh!