1#import bevy_pbr::meshlet_bindings::{
2    InstancedOffset,
3    constants,
4    meshlet_view_instance_visibility,
5    meshlet_instance_aabbs,
6    meshlet_instance_bvh_root_nodes,
7    meshlet_bvh_cull_count_write,
8    meshlet_bvh_cull_dispatch,
9    meshlet_bvh_cull_queue,
10    meshlet_second_pass_instance_count,
11    meshlet_second_pass_instance_dispatch,
12    meshlet_second_pass_instance_candidates,
13}
14#import bevy_pbr::meshlet_cull_shared::{
15    aabb_in_frustum,
16    should_occlusion_cull_aabb,
17}
18
19fn instance_count() -> u32 {
20#ifdef MESHLET_FIRST_CULLING_PASS
21    return constants.scene_instance_count;
22#else
23    return meshlet_second_pass_instance_count;
24#endif
25}
26
27fn map_instance_id(id: u32) -> u32 {
28#ifdef MESHLET_FIRST_CULLING_PASS
29    return id;
30#else
31    return meshlet_second_pass_instance_candidates[id];
32#endif
33}
34
35fn should_cull_instance(instance_id: u32) -> bool {
36    let index = instance_id >> 5u;
37    let bit_offset = instance_id & 31u;
38    let packed_visibility = meshlet_view_instance_visibility[index];
39    return bool(extractBits(packed_visibility, bit_offset, 1u));
40}
41
42@compute
43@workgroup_size(128, 1, 1) // 1 instance per thread
44fn cull_instances(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {
45    // Calculate the instance ID for this thread
46    let dispatch_id = global_invocation_id.x;
47    if dispatch_id >= instance_count() { return; }
48
49    let instance_id = map_instance_id(dispatch_id);
50    let aabb = meshlet_instance_aabbs[instance_id];
51
52    // Visibility and frustum cull, but only in the first pass
53#ifdef MESHLET_FIRST_CULLING_PASS
54    if should_cull_instance(instance_id) || !aabb_in_frustum(aabb, instance_id) { return; }
55#endif
56
57    // If we pass, try occlusion culling
58    // If this instance was occluded, push it to the second pass to check against this frame's HZB
59    if should_occlusion_cull_aabb(aabb, instance_id) {
60#ifdef MESHLET_FIRST_CULLING_PASS
61        let id = atomicAdd(&meshlet_second_pass_instance_count, 1u);
62        meshlet_second_pass_instance_candidates[id] = instance_id;
63        if ((id & 127u) == 0u) {
64            atomicAdd(&meshlet_second_pass_instance_dispatch.x, 1u);
65        }
66#endif
67        return;
68    }
69
70    // If we pass, push the instance's root node to BVH cull
71    let root_node = meshlet_instance_bvh_root_nodes[instance_id];
72    let id = atomicAdd(&meshlet_bvh_cull_count_write, 1u);
73    meshlet_bvh_cull_queue[id] = InstancedOffset(instance_id, root_node);
74    if ((id & 15u) == 0u) {
75        atomicAdd(&meshlet_bvh_cull_dispatch.x, 1u);
76    }
77}