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