#import nightshade_renderer::meshlet_data::{Meshlet, meshlet_payload, get_meshlet_triangle_count}
#import nightshade_renderer::meshlet_streams::{raster_clusters, meshlets, meshlet_instances, get_meshlet_vertex_id, get_meshlet_vertex_position}
/// Rasterizes small clusters in compute, without the fixed function rasterizer.
///
/// A hardware rasterizer shades in two by two quads so it can take derivatives,
/// so a triangle covering a single pixel occupies four lanes and wastes three. A
/// level of detail cut drives triangles down to about a pixel, so at the size
/// virtualized geometry draws, much of that hardware goes to quads nothing asked
/// for.
///
/// This pass owns the small end. One workgroup takes one cluster, one thread
/// takes one triangle, and each thread walks only the pixels its own triangle
/// covers. Depth rides in the high half of a sixty four bit value with the
/// cluster and triangle in the low half, so a single atomic max is both the
/// depth test and the write: the largest value wins, and under a reversed depth
/// the largest is the nearest. No depth attachment, no blending, no quads.
///
/// Clusters too large for this, and any crossing the near plane and needing
/// clipping, go to the hardware path instead.
struct MeshletView {
clip_from_world: mat4x4<f32>,
camera_position: vec4<f32>,
screen_size: vec4<f32>,
counts: vec4<u32>,
}
/// Only the length is read here. The counts in front of it belong to the
/// dispatch that launched this.
struct DispatchIndirect {
workgroup_count_x: u32,
workgroup_count_y: u32,
workgroup_count_z: u32,
total_clusters: u32,
software_clusters: u32,
}
@group(0) @binding(0) var<uniform> view: MeshletView;
@group(0) @binding(1) var visibility_buffer: texture_storage_2d<r64uint, atomic>;
@group(0) @binding(2) var<storage, read> dispatch_args: DispatchIndirect;
/// Twice the signed area of the triangle `(a, b, point)` in pixels. Its sign is
/// the winding, and it is the denominator that turns edge functions into
/// barycentrics.
fn edge_function(a: vec2<f32>, b: vec2<f32>, point: vec2<f32>) -> f32 {
return (b.x - a.x) * (point.y - a.y) - (b.y - a.y) * (point.x - a.x);
}
/// Writes a covered pixel, if it is nearer than whatever is already there.
///
/// Depth occupies the high half so comparing the whole value compares depth
/// first. The depth is reversed, so nearer is larger and the max wins.
///
/// A depth of exactly zero is the far plane and is skipped, which keeps every
/// value this writes non-zero, so a cleared texel means nothing rasterized here
/// without a sentinel colliding with real geometry.
fn write_visibility(pixel: vec2<i32>, depth: f32, payload: u32) {
let depth_bits = bitcast<u32>(depth);
if depth_bits == 0u {
return;
}
let packed = (u64(depth_bits) << 32u) | u64(payload);
textureAtomicMax(visibility_buffer, pixel, packed);
}
@compute @workgroup_size(#{MESHLET_MAX_TRIANGLES})
fn software_raster_main(
@builtin(workgroup_id) workgroup_id: vec3<u32>,
@builtin(num_workgroups) num_workgroups: vec3<u32>,
@builtin(local_invocation_id) local_invocation_id: vec3<u32>,
) {
// The list is folded across two dimensions because one may not be dispatched
// far enough to hold it, so the row width is where the fold happened and the
// last row runs past the end.
let cluster_id = workgroup_id.y * num_workgroups.x + workgroup_id.x;
if cluster_id >= dispatch_args.software_clusters {
return;
}
let cluster = raster_clusters[cluster_id];
var meshlet = meshlets[cluster.offset];
let triangle_id = local_invocation_id.x;
if triangle_id >= get_meshlet_triangle_count(&meshlet) {
return;
}
let instance = meshlet_instances[cluster.instance_id];
let clip_from_local = view.clip_from_world * instance.world_from_local;
let index_id = meshlet.start_index_id + triangle_id * 3u;
var clip_positions: array<vec4<f32>, 3>;
for (var corner = 0u; corner < 3u; corner++) {
let vertex_id = get_meshlet_vertex_id(index_id + corner);
clip_positions[corner] = clip_from_local
* vec4<f32>(get_meshlet_vertex_position(&meshlet, vertex_id), 1.0);
}
// Nothing here clips, so a vertex at or behind the eye would project through
// infinity. The cull keeps those clusters away from this pass; this guards a
// triangle that slips through anyway.
if clip_positions[0].w <= 0.0 || clip_positions[1].w <= 0.0 || clip_positions[2].w <= 0.0 {
return;
}
// Screen space in pixels, carrying the reversed depth along unchanged: depth
// over w is affine across the screen, so plain barycentrics interpolate it
// exactly and no perspective correction belongs here.
var screen_positions: array<vec3<f32>, 3>;
for (var corner = 0u; corner < 3u; corner++) {
let ndc = clip_positions[corner].xyz / clip_positions[corner].w;
screen_positions[corner] = vec3<f32>(
(ndc.x * 0.5 + 0.5) * view.screen_size.x,
(0.5 - ndc.y * 0.5) * view.screen_size.y,
ndc.z,
);
}
let area = edge_function(
screen_positions[0].xy,
screen_positions[1].xy,
screen_positions[2].xy,
);
if abs(area) < 1.0e-9 {
return;
}
let minimum = min(
screen_positions[0].xy,
min(screen_positions[1].xy, screen_positions[2].xy),
);
let maximum = max(
screen_positions[0].xy,
max(screen_positions[1].xy, screen_positions[2].xy),
);
let bounds_minimum = vec2<i32>(max(floor(minimum), vec2<f32>(0.0)));
let bounds_maximum = vec2<i32>(min(ceil(maximum), view.screen_size.xy - vec2<f32>(1.0)));
let payload = meshlet_payload(cluster_id, triangle_id);
// Dividing by the signed area rather than its magnitude normalizes both
// windings to positive weights inside the triangle, which is what the
// hardware path does by rasterizing with no face culling at all. Taking the
// magnitude here would drop every back face and leave the two paths
// disagreeing about the same mesh.
let inverse_area = 1.0 / area;
for (var y = bounds_minimum.y; y <= bounds_maximum.y; y++) {
for (var x = bounds_minimum.x; x <= bounds_maximum.x; x++) {
let sample = vec2<f32>(f32(x) + 0.5, f32(y) + 0.5);
let weight_a = edge_function(
screen_positions[1].xy,
screen_positions[2].xy,
sample,
) * inverse_area;
let weight_b = edge_function(
screen_positions[2].xy,
screen_positions[0].xy,
sample,
) * inverse_area;
let weight_c = 1.0 - weight_a - weight_b;
if weight_a < 0.0 || weight_b < 0.0 || weight_c < 0.0 {
continue;
}
let depth = screen_positions[0].z * weight_a
+ screen_positions[1].z * weight_b
+ screen_positions[2].z * weight_c;
write_visibility(vec2<i32>(x, y), depth, payload);
}
}
}
/// Zeroes the visibility buffer. A storage texture takes no attachment clear, so
/// the clear is a pass of its own.
@compute @workgroup_size(8, 8)
fn clear_main(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {
let pixel = vec2<i32>(global_invocation_id.xy);
if pixel.x >= i32(view.screen_size.x) || pixel.y >= i32(view.screen_size.y) {
return;
}
textureStore(visibility_buffer, pixel, vec4<u64>(u64(0u)));
}