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 bit_offset = instance_id >> 5u;
37    let packed_visibility = meshlet_view_instance_visibility[instance_id & 31u];
38    return bool(extractBits(packed_visibility, bit_offset, 1u));
39}
40
41@compute
42@workgroup_size(128, 1, 1) // 1 instance per thread
43fn cull_instances(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {
44    // Calculate the instance ID for this thread
45    let dispatch_id = global_invocation_id.x;
46    if dispatch_id >= instance_count() { return; }
47
48    let instance_id = map_instance_id(dispatch_id);
49    let aabb = meshlet_instance_aabbs[instance_id];
50
51    // Visibility and frustum cull, but only in the first pass
52#ifdef MESHLET_FIRST_CULLING_PASS
53    if should_cull_instance(instance_id) || !aabb_in_frustum(aabb, instance_id) { return; }
54#endif
55
56    // If we pass, try occlusion culling
57    // If this instance was occluded, push it to the second pass to check against this frame's HZB
58    if should_occlusion_cull_aabb(aabb, instance_id) {
59#ifdef MESHLET_FIRST_CULLING_PASS
60        let id = atomicAdd(&meshlet_second_pass_instance_count, 1u);
61        meshlet_second_pass_instance_candidates[id] = instance_id;
62        if ((id & 127u) == 0u) {
63            atomicAdd(&meshlet_second_pass_instance_dispatch.x, 1u);
64        }
65#endif
66        return;
67    }
68
69    // If we pass, push the instance's root node to BVH cull
70    let root_node = meshlet_instance_bvh_root_nodes[instance_id];
71    let id = atomicAdd(&meshlet_bvh_cull_count_write, 1u);
72    meshlet_bvh_cull_queue[id] = InstancedOffset(instance_id, root_node);
73    if ((id & 15u) == 0u) {
74        atomicAdd(&meshlet_bvh_cull_dispatch.x, 1u);
75    }
76}