#import nightshade_renderer::cull_common::sphere_in_frustum
#import nightshade_renderer::meshlet_data::{MeshletInstance, InstancedOffset}
/// Picks the clusters that make up one cut through every instance's level of
/// detail hierarchy, and builds the raster's draw from them.
///
/// One thread owns one instance and walks that instance's bvh alone, so the
/// cpu never iterates the scene: it uploads instances as they change and issues
/// a single indirect draw whose instance count this pass wrote.
/// Depth times the eight children a node can hold. The bake keeps trees
/// shallow, so this holds a full walk without spilling.
const CULL_STACK_CAPACITY: u32 = 64u;
struct MeshletAabbErrorOffset {
center: vec3<f32>,
error: f32,
half_extent: vec3<f32>,
child_offset: u32,
}
struct MeshletBoundingSphere {
center: vec3<f32>,
radius: f32,
}
/// Mirrors the baked layout exactly: `child_counts` is eight bytes, which wgsl
/// can only address as two words, so a slot's count is extracted rather than
/// indexed.
struct BvhNode {
aabbs: array<MeshletAabbErrorOffset, 8>,
lod_bounds: array<MeshletBoundingSphere, 8>,
child_counts: array<u32, 2>,
padding: array<u32, 2>,
}
struct MeshletCullData {
aabb: MeshletAabbErrorOffset,
lod_group_sphere: MeshletBoundingSphere,
}
struct DrawIndirect {
vertex_count: u32,
instance_count: atomic<u32>,
first_vertex: u32,
first_instance: u32,
}
/// A dispatch reads only the three counts. The rest rides behind them.
struct DispatchIndirect {
workgroup_count_x: u32,
workgroup_count_y: u32,
workgroup_count_z: u32,
total_clusters: atomic<u32>,
software_clusters: atomic<u32>,
}
/// `params` is (pixels per unit at unit depth, near plane, orthographic pixels
/// per unit or zero when perspective, the error in pixels a cluster may
/// introduce). `counts` is (instances, capacity of the cluster list, the widest
/// a cluster may appear and still rasterize in compute, non-zero when a compute
/// rasterizer exists at all). `limits.x` is the most workgroups the device will
/// dispatch along one dimension; `limits.y` is non-zero when the occlusion test
/// has a depth pyramid to read; `limits.z` is that pyramid's mip count.
/// `occluder_from_world` and `occluder_screen_size` project a cluster's box into
/// the pyramid, which is the previous frame's depth in that camera's screen.
struct MeshletCullView {
frustum_planes: array<vec4<f32>, 6>,
occluder_from_world: mat4x4<f32>,
camera_position: vec4<f32>,
params: vec4<f32>,
counts: vec4<u32>,
limits: vec4<u32>,
occluder_screen_size: vec4<f32>,
}
@group(0) @binding(0) var<storage, read> meshlet_instances: array<MeshletInstance>;
@group(0) @binding(1) var<storage, read> meshlet_bvh_nodes: array<BvhNode>;
@group(0) @binding(2) var<storage, read> meshlet_cull_data: array<MeshletCullData>;
@group(0) @binding(3) var<storage, read_write> raster_clusters: array<InstancedOffset>;
@group(0) @binding(4) var<storage, read_write> draw_args: DrawIndirect;
@group(0) @binding(5) var<uniform> view: MeshletCullView;
@group(0) @binding(6) var<storage, read_write> dispatch_args: DispatchIndirect;
@group(0) @binding(7) var hiz_pyramid: texture_2d<f32>;
fn get_child_count(node: ptr<function, BvhNode>, slot: u32) -> u32 {
let packed = (*node).child_counts[slot >> 2u];
return extractBits(packed, (slot & 3u) * 8u, 8u);
}
fn transform_scale(world_from_local: mat4x4<f32>) -> f32 {
return max(
length(world_from_local[0].xyz),
max(length(world_from_local[1].xyz), length(world_from_local[2].xyz)),
);
}
/// Whether `error`, in the mesh's local units inside `sphere`, stays under a
/// pixel once projected. Distance is held at the near plane: nothing closer is
/// drawn, so letting it fall towards zero would report an unbounded error and
/// drive the cut to its finest level for a cluster the camera sits inside.
///
/// The error is carried into world space before it is projected. A bake
/// measures its error in local units, and the distance it is divided by is a
/// world length, so an instance scaled up ten times introduces ten times the
/// error at the same distance and must reach for a finer level sooner.
/// Omitting the scale would leave the two sides of that division in different
/// spaces and pick geometry too coarse for anything scaled up. This is why the
/// scale applies to both projections rather than only the orthographic one.
fn error_is_imperceptible(
sphere: MeshletBoundingSphere,
error: f32,
world_from_local: mat4x4<f32>,
scale: f32,
) -> bool {
let world_error = error * scale;
let orthographic_pixels_per_unit = view.params.z;
if orthographic_pixels_per_unit > 0.0 {
return world_error * orthographic_pixels_per_unit < view.params.w;
}
let center = (world_from_local * vec4<f32>(sphere.center, 1.0)).xyz;
let distance = length(center - view.camera_position.xyz) - sphere.radius * scale;
let projected = world_error * view.params.x / max(distance, view.params.y);
return projected < view.params.w;
}
fn sphere_is_visible(
sphere: MeshletBoundingSphere,
world_from_local: mat4x4<f32>,
scale: f32,
) -> bool {
let center = (world_from_local * vec4<f32>(sphere.center, 1.0)).xyz;
return sphere_in_frustum(view.frustum_planes, center, sphere.radius * scale);
}
/// How wide `sphere` appears, in pixels, and whether it reaches nearer than the
/// near plane.
///
/// A compute rasterizer has no clipping hardware, so a cluster that crosses the
/// near plane cannot go to one at any size: its vertices project through
/// infinity. That case is rare and enormous, which is the hardware path's to
/// begin with.
struct ClusterScreenExtent {
pixels: f32,
crosses_near_plane: bool,
}
fn cluster_screen_extent(
sphere: MeshletBoundingSphere,
world_from_local: mat4x4<f32>,
scale: f32,
) -> ClusterScreenExtent {
let radius = sphere.radius * scale;
let orthographic_pixels_per_unit = view.params.z;
if orthographic_pixels_per_unit > 0.0 {
return ClusterScreenExtent(radius * 2.0 * orthographic_pixels_per_unit, false);
}
let center = (world_from_local * vec4<f32>(sphere.center, 1.0)).xyz;
let distance = length(center - view.camera_position.xyz);
let near = view.params.y;
return ClusterScreenExtent(
radius * 2.0 * view.params.x / max(distance - radius, near),
distance - radius < near,
);
}
/// Claims a slot for a cluster in whichever list rasterizes it.
///
/// Both lists live in one array and grow towards each other, software from the
/// front and hardware from the back, so a slot is only claimed once the shared
/// counter has admitted it under the capacity. That ordering is what keeps the
/// two ends apart: admitted clusters never outnumber the slots, so the front
/// index and the back index cannot reach the same one. Deciding from each list's
/// own length instead would let both ends read a safe-looking length in the same
/// instant and write the same slot.
///
/// The shared counter keeps counting past capacity on purpose, so it holds the
/// number of clusters the cut kept rather than the number that fit.
fn emit_cluster(instance_id: u32, meshlet_index: u32, extent: ClusterScreenExtent) {
let admitted = atomicAdd(&dispatch_args.total_clusters, 1u);
if admitted >= view.counts.y {
return;
}
let rasterizes_in_compute = view.counts.w != 0u
&& !extent.crosses_near_plane
&& extent.pixels <= f32(view.counts.z);
if rasterizes_in_compute {
let slot = atomicAdd(&dispatch_args.software_clusters, 1u);
raster_clusters[slot] = InstancedOffset(instance_id, meshlet_index);
return;
}
let slot = atomicAdd(&draw_args.instance_count, 1u);
raster_clusters[view.counts.y - 1u - slot] = InstancedOffset(instance_id, meshlet_index);
}
/// Whether a local-space box, placed by `world_from_local`, is hidden behind the
/// depth pyramid. The pyramid is the previous frame's depth in the occluder
/// camera's screen: for a scene that barely moves between frames, which is what
/// virtualized geometry is for, it predicts this frame closely, and the cost of
/// being a frame stale is a cluster drawn a frame late rather than a wrong one.
///
/// The box's eight corners project to a screen rectangle and a nearest depth.
/// The pyramid holds the farthest surface over each region, so a mip chosen to
/// cover the rectangle in about one texel answers whether even that farthest
/// occluder is nearer than the box: if so, nothing in the rectangle leaves the
/// box visible and it is dropped. Reverse-Z runs throughout, so nearer is the
/// larger value and the pyramid reduces by minimum. A box crossing the near
/// plane cannot be projected and is kept.
fn box_is_occluded(
local_center: vec3<f32>,
local_half_extent: vec3<f32>,
world_from_local: mat4x4<f32>,
) -> bool {
if view.limits.y == 0u {
return false;
}
var screen_min = vec2<f32>(1.0);
var screen_max = vec2<f32>(0.0);
var nearest_z = 0.0;
for (var corner = 0u; corner < 8u; corner++) {
let corner_sign = vec3<f32>(
select(-1.0, 1.0, (corner & 1u) != 0u),
select(-1.0, 1.0, (corner & 2u) != 0u),
select(-1.0, 1.0, (corner & 4u) != 0u),
);
let world = (world_from_local * vec4<f32>(local_center + corner_sign * local_half_extent, 1.0)).xyz;
let clip = view.occluder_from_world * vec4<f32>(world, 1.0);
if clip.w <= 0.0 {
return false;
}
let ndc = clip.xyz / clip.w;
let uv = vec2<f32>(ndc.x * 0.5 + 0.5, 1.0 - (ndc.y * 0.5 + 0.5));
screen_min = min(screen_min, uv);
screen_max = max(screen_max, uv);
nearest_z = max(nearest_z, ndc.z);
}
let pixel_pad = vec2<f32>(1.0) / view.occluder_screen_size.xy;
let padded_min = clamp(screen_min - pixel_pad, vec2<f32>(0.0), vec2<f32>(1.0));
let padded_max = clamp(screen_max + pixel_pad, vec2<f32>(0.0), vec2<f32>(1.0));
if padded_min.x >= padded_max.x || padded_min.y >= padded_max.y {
return false;
}
let rect_size = max(
(padded_max.x - padded_min.x) * view.occluder_screen_size.x,
(padded_max.y - padded_min.y) * view.occluder_screen_size.y,
);
// One level finer than the mip whose texel matches the footprint, so the
// box spans about two texels an axis rather than one. A single matching
// texel is reduced over the box plus as much again around it, and that
// skirt reaches into a neighbouring donut's hole and reads the far surface
// there as visible, leaving a hidden cluster drawn. Two half-sized texels
// keep the minimum inside the box, which is what actually occludes it.
let mip = i32(clamp(ceil(log2(max(rect_size, 1.0))) - 1.0, 0.0, f32(view.limits.z) - 1.0));
let mip_size = vec2<f32>(textureDimensions(hiz_pyramid, mip));
let min_texel = clamp(vec2<i32>(floor(padded_min * mip_size)), vec2<i32>(0), vec2<i32>(mip_size) - vec2<i32>(1));
let max_texel = clamp(vec2<i32>(floor(padded_max * mip_size)), vec2<i32>(0), vec2<i32>(mip_size) - vec2<i32>(1));
var farthest_occluder = 1.0;
for (var y = min_texel.y; y <= max_texel.y; y++) {
for (var x = min_texel.x; x <= max_texel.x; x++) {
farthest_occluder = min(farthest_occluder, textureLoad(hiz_pyramid, vec2<i32>(x, y), mip).r);
}
}
// Relative, with no fixed floor: under reverse-Z with an infinite far plane
// the whole scene's depth lives in a sliver near zero, so a constant bias of
// a few thousandths is larger than the depths themselves and pushes the
// threshold negative, culling nothing. A fraction of the occluder's own
// depth scales with wherever the scene sits.
let bias = farthest_occluder * 0.02;
return nearest_z < farthest_occluder - bias;
}
/// Every slot carries the error that using its parent would introduce, and that
/// error only grows towards the root, so a slot already under a pixel means the
/// whole subtree is better served by coarser geometry drawn elsewhere: prune
/// it. Where the parent's error is too large the group is a candidate, and a
/// meshlet in it draws only if its own error is acceptable. Those two tests
/// bracket exactly one level per region, which is what keeps levels from
/// overlapping and from leaving cracks.
///
/// A slot hidden behind the depth pyramid is pruned before either, which drops
/// whole subtrees the scene already covers. The slot's box bounds everything
/// under it, so one test at the slot removes the group or the subtree alike.
@compute @workgroup_size(64)
fn cull_main(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {
let instance_id = global_invocation_id.x;
if instance_id >= view.counts.x {
return;
}
let instance = meshlet_instances[instance_id];
let world_from_local = instance.world_from_local;
let scale = transform_scale(world_from_local);
var stack: array<u32, CULL_STACK_CAPACITY>;
var stack_size = 1u;
stack[0] = instance.root_bvh_node_index;
while stack_size > 0u {
stack_size -= 1u;
var node = meshlet_bvh_nodes[stack[stack_size]];
for (var slot = 0u; slot < 8u; slot++) {
let child_count = get_child_count(&node, slot);
if child_count == 0u {
break;
}
let lod_sphere = node.lod_bounds[slot];
if error_is_imperceptible(lod_sphere, node.aabbs[slot].error, world_from_local, scale) {
continue;
}
if !sphere_is_visible(lod_sphere, world_from_local, scale) {
continue;
}
if box_is_occluded(node.aabbs[slot].center, node.aabbs[slot].half_extent, world_from_local) {
continue;
}
let child_offset = node.aabbs[slot].child_offset;
if child_count == 255u {
if stack_size < CULL_STACK_CAPACITY {
stack[stack_size] = child_offset;
stack_size += 1u;
}
continue;
}
for (var meshlet_offset = 0u; meshlet_offset < child_count; meshlet_offset++) {
let meshlet_index = child_offset + meshlet_offset;
let cull_data = meshlet_cull_data[meshlet_index];
if !error_is_imperceptible(
cull_data.lod_group_sphere,
cull_data.aabb.error,
world_from_local,
scale,
) {
continue;
}
if !sphere_is_visible(cull_data.lod_group_sphere, world_from_local, scale) {
continue;
}
// Tested per cluster, not only at the group above: a group box
// bounds several clusters and is too large to ever fall wholly
// behind an occluder, so occlusion has to reach the single
// cluster's box to cull anything a nearer surface hides.
if box_is_occluded(cull_data.aabb.center, cull_data.aabb.half_extent, world_from_local) {
continue;
}
// Measured from the cluster's own box rather than the sphere
// that bounds its whole level of detail group: the group is
// several clusters wide, so it would report every one of them
// as far too large for the compute path.
emit_cluster(
instance_id,
meshlet_index,
cluster_screen_extent(
MeshletBoundingSphere(
cull_data.aabb.center,
length(cull_data.aabb.half_extent),
),
world_from_local,
scale,
),
);
}
}
}
}
/// Turns the software list's length into a dispatch that respects the bound on
/// workgroups per dimension.
///
/// One workgroup rasterizes one cluster, and the list is far longer than a
/// single dimension may be dispatched: a scene of any size passes it easily. A
/// dispatch that asks for more than the device allows is undefined rather than
/// clamped, and an indirect one is never checked, because the count only exists
/// on the gpu by the time it is read. So the list is folded across two
/// dimensions, and the rasterizer recovers the index from both. The last row is
/// partial, which is why the length travels separately for it to test against.
@compute @workgroup_size(1)
fn software_dispatch_main() {
let clusters = atomicLoad(&dispatch_args.software_clusters);
let stride = max(view.limits.x, 1u);
dispatch_args.workgroup_count_x = min(clusters, stride);
dispatch_args.workgroup_count_y = (clusters + stride - 1u) / stride;
dispatch_args.workgroup_count_z = 1u;
}