From 175e1462289c217ec57a3d3b9894a7c513890fa4 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Thu, 6 Jun 2024 10:10:07 -0700 Subject: [PATCH] Misc meshlet changes (#13705) * Rename cull_meshlets -> cull_clusters * Rename meshlet_visible -> cluster_visible * Add an if statement around meshlet_second_pass_candidates writes, maybe a small bit of performance. --- .../{cull_meshlets.wgsl => cull_clusters.wgsl} | 16 +++++++++------- crates/bevy_pbr/src/meshlet/mod.rs | 2 +- crates/bevy_pbr/src/meshlet/pipelines.rs | 4 ++-- 3 files changed, 12 insertions(+), 10 deletions(-) rename crates/bevy_pbr/src/meshlet/{cull_meshlets.wgsl => cull_clusters.wgsl} (96%) diff --git a/crates/bevy_pbr/src/meshlet/cull_meshlets.wgsl b/crates/bevy_pbr/src/meshlet/cull_clusters.wgsl similarity index 96% rename from crates/bevy_pbr/src/meshlet/cull_meshlets.wgsl rename to crates/bevy_pbr/src/meshlet/cull_clusters.wgsl index e25532b319..373b1e411d 100644 --- a/crates/bevy_pbr/src/meshlet/cull_meshlets.wgsl +++ b/crates/bevy_pbr/src/meshlet/cull_clusters.wgsl @@ -22,7 +22,7 @@ @compute @workgroup_size(128, 1, 1) // 128 threads per workgroup, 1 cluster per thread -fn cull_meshlets( +fn cull_clusters( @builtin(workgroup_id) workgroup_id: vec3, @builtin(num_workgroups) num_workgroups: vec3, @builtin(local_invocation_id) local_invocation_id: vec3, @@ -100,25 +100,27 @@ fn cull_meshlets( let occluder_depth = min(min(depth_quad_a, depth_quad_b), min(depth_quad_c, depth_quad_d)); // Check whether or not the cluster would be occluded if drawn - var meshlet_visible: bool; + var cluster_visible: bool; if view.clip_from_view[3][3] == 1.0 { // Orthographic let sphere_depth = view.clip_from_view[3][2] + (culling_bounding_sphere_center_view_space.z + culling_bounding_sphere_radius) * view.clip_from_view[2][2]; - meshlet_visible = sphere_depth >= occluder_depth; + cluster_visible = sphere_depth >= occluder_depth; } else { // Perspective let sphere_depth = -view.clip_from_view[3][2] / (culling_bounding_sphere_center_view_space.z + culling_bounding_sphere_radius); - meshlet_visible = sphere_depth >= occluder_depth; + cluster_visible = sphere_depth >= occluder_depth; } // Write if the cluster should be occlusion tested in the second pass #ifdef MESHLET_FIRST_CULLING_PASS - let second_pass_candidate = u32(!meshlet_visible) << cluster_id % 32u; - atomicOr(&meshlet_second_pass_candidates[cluster_id / 32u], second_pass_candidate); + if !cluster_visible { + let bit = 1u << cluster_id % 32u; + atomicOr(&meshlet_second_pass_candidates[cluster_id / 32u], bit); + } #endif // Append a list of this cluster's triangles to draw if not culled - if meshlet_visible { + if cluster_visible { let meshlet_triangle_count = meshlets[meshlet_id].triangle_count; let buffer_start = atomicAdd(&draw_indirect_args.vertex_count, meshlet_triangle_count * 3u) / 3u; let cluster_id_packed = cluster_id << 6u; diff --git a/crates/bevy_pbr/src/meshlet/mod.rs b/crates/bevy_pbr/src/meshlet/mod.rs index 69255f9084..fd73e35c4a 100644 --- a/crates/bevy_pbr/src/meshlet/mod.rs +++ b/crates/bevy_pbr/src/meshlet/mod.rs @@ -136,7 +136,7 @@ impl Plugin for MeshletPlugin { load_internal_asset!( app, MESHLET_CULLING_SHADER_HANDLE, - "cull_meshlets.wgsl", + "cull_clusters.wgsl", Shader::from_wgsl ); load_internal_asset!( diff --git a/crates/bevy_pbr/src/meshlet/pipelines.rs b/crates/bevy_pbr/src/meshlet/pipelines.rs index 2feaf9518b..0112dcbb67 100644 --- a/crates/bevy_pbr/src/meshlet/pipelines.rs +++ b/crates/bevy_pbr/src/meshlet/pipelines.rs @@ -67,7 +67,7 @@ impl FromWorld for MeshletPipelines { "MESHLET_CULLING_PASS".into(), "MESHLET_FIRST_CULLING_PASS".into(), ], - entry_point: "cull_meshlets".into(), + entry_point: "cull_clusters".into(), }), cull_second: pipeline_cache.queue_compute_pipeline(ComputePipelineDescriptor { @@ -79,7 +79,7 @@ impl FromWorld for MeshletPipelines { "MESHLET_CULLING_PASS".into(), "MESHLET_SECOND_CULLING_PASS".into(), ], - entry_point: "cull_meshlets".into(), + entry_point: "cull_clusters".into(), }), downsample_depth_first: pipeline_cache.queue_compute_pipeline(