1import package::meshlet::bindings::{
2    InstancedOffset,
3    get_aabb,
4    get_aabb_error,
5    constants,
6    view,
7    meshlet_instance_uniforms,
8    meshlet_cull_data,
9    meshlet_software_raster_indirect_args,
10    meshlet_hardware_raster_indirect_args,
11    meshlet_previous_raster_counts,
12    meshlet_raster_clusters,
13    meshlet_meshlet_cull_count_read,
14    meshlet_meshlet_cull_count_write,
15    meshlet_meshlet_cull_dispatch,
16    meshlet_meshlet_cull_queue,
17};
18import package::meshlet::cull_shared::{
19    ScreenAabb,
20    project_aabb,
21    lod_error_is_imperceptible,
22    aabb_in_frustum,
23    should_occlusion_cull_aabb,
24};
25import bevy_render::maths::affine3_to_square;
26
27@compute
28@workgroup_size(128, 1, 1) // 1 cluster per thread
29fn cull_clusters(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {
30    if global_invocation_id.x >= meshlet_meshlet_cull_count_read { return; }
31
32    @if(MESHLET_FIRST_CULLING_PASS)
33    let meshlet_id = global_invocation_id.x;
34    @else
35    let meshlet_id = constants.rightmost_slot - global_invocation_id.x;
36    let instanced_offset = meshlet_meshlet_cull_queue[meshlet_id];
37    let instance_id = instanced_offset.instance_id;
38    let cull_data = &meshlet_cull_data[instanced_offset.offset];
39    var aabb_error_offset = (*cull_data).aabb;
40    let aabb = get_aabb(&aabb_error_offset);
41    let error = get_aabb_error(&aabb_error_offset);
42    let lod_sphere = (*cull_data).lod_group_sphere;
43
44    let is_imperceptible = lod_error_is_imperceptible(lod_sphere, error, instance_id);
45    // Error and frustum cull, in both passes
46    if !is_imperceptible || !aabb_in_frustum(aabb, instance_id) { return; }
47
48    // If we pass, try occlusion culling
49    // If this node was occluded, push it's children to the second pass to check against this frame's HZB
50    if should_occlusion_cull_aabb(aabb, instance_id) {
51@if(MESHLET_FIRST_CULLING_PASS) {
52        let id = atomicAdd(&meshlet_meshlet_cull_count_write, 1u);
53        let value = InstancedOffset(instance_id, instanced_offset.offset);
54        meshlet_meshlet_cull_queue[constants.rightmost_slot - id] = value;
55        if ((id & 127u) == 0) {
56            atomicAdd(&meshlet_meshlet_cull_dispatch.x, 1u);
57        }
58}
59        return;
60    }
61
62    // If we pass, rasterize the meshlet
63    // Check how big the cluster is in screen space
64    let world_from_local = affine3_to_square(meshlet_instance_uniforms[instance_id].world_from_local);
65    let clip_from_local  = view.clip_from_world * world_from_local;
66    let projection = view.clip_from_world;
67    var near: f32;
68    if projection[3][3] == 1.0 {
69        near = projection[3][2] / projection[2][2];
70    } else {
71        near = projection[3][2];
72    }
73    var screen_aabb = ScreenAabb(vec3<f32>(0.0), vec3<f32>(0.0));
74    var sw_raster = project_aabb(clip_from_local, near, aabb, &screen_aabb);
75    if sw_raster {
76        let aabb_size = (screen_aabb.max.xy - screen_aabb.min.xy) * view.viewport.zw;
77        sw_raster = all(aabb_size <= vec2<f32>(64.0));
78    }
79
80    var buffer_slot: u32;
81    if sw_raster {
82        // Append this cluster to the list for software rasterization
83        buffer_slot = atomicAdd(&meshlet_software_raster_indirect_args.x, 1u);
84        buffer_slot += meshlet_previous_raster_counts[0];
85    } else {
86        // Append this cluster to the list for hardware rasterization
87        buffer_slot = atomicAdd(&meshlet_hardware_raster_indirect_args.instance_count, 1u);
88        buffer_slot += meshlet_previous_raster_counts[1];
89        buffer_slot = constants.rightmost_slot - buffer_slot;
90    }
91    meshlet_raster_clusters[buffer_slot] = InstancedOffset(instance_id, instanced_offset.offset);
92}