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.
This commit is contained in:
JMS55 2024-06-06 10:10:07 -07:00 committed by GitHub
parent 31be32ff10
commit 175e146228
No known key found for this signature in database
GPG key ID: B5690EEEBB952194
3 changed files with 12 additions and 10 deletions

View file

@ -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<u32>,
@builtin(num_workgroups) num_workgroups: vec3<u32>,
@builtin(local_invocation_id) local_invocation_id: vec3<u32>,
@ -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;

View file

@ -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!(

View file

@ -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(