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