// Light culling for the clustered forward pass.
//
// The frustum is cut into a grid of cells -- froxels, because they are frusta
// rather than boxes -- and every cell is told which lights reach it. A
// fragment then answers only to the lights in the one cell it sits in, which
// is what makes a scene with two hundred lamps cost about what a scene with
// four does.
//
// The slices go back exponentially, not evenly: perspective already spends
// most of the depth buffer on what is near, and even slices would put twenty
// of the twenty-four cells in the far distance where a local light is a
// pinprick. `z = near * (far/near)^(k/slices)` gives every slice the same
// ratio of depths, which is the same reasoning as a musical scale.
struct Clusters {
// Screen back to view space, for the corners of a cell.
inv_proj: mat4x4<f32>,
// x, y: the frame in pixels. z, w: near and far.
screen: vec4<f32>,
// x, y, z: how many cells each way. w: how many lights there are.
grid: vec4<u32>,
// x, y: scale and bias taking log(view depth) to a slice.
slice: vec4<f32>,
};
struct LocalLight {
position_range: vec4<f32>,
color: vec4<f32>,
direction_outer: vec4<f32>,
cone: vec4<f32>,
};
@group(0) @binding(0) var<uniform> clusters: Clusters;
@group(0) @binding(1) var<storage, read> lights: array<LocalLight>;
@group(0) @binding(2) var<storage, read_write> cluster_counts: array<u32>;
@group(0) @binding(3) var<storage, read_write> cluster_lights: array<u32>;
// The view matrix, to put a light's world position into the space the cells
// are built in.
@group(0) @binding(4) var<uniform> view: mat4x4<f32>;
const MAX_LIGHTS_PER_CLUSTER: u32 = 64u;
// A point on the near plane, in view space, under a pixel.
fn view_ray(pixel: vec2<f32>) -> vec3<f32> {
let ndc = vec4<f32>(
2.0 * pixel.x / clusters.screen.x - 1.0,
1.0 - 2.0 * pixel.y / clusters.screen.y,
1.0,
1.0,
);
let v = clusters.inv_proj * ndc;
return v.xyz / v.w;
}
// Where a ray from the eye through `p` crosses the plane at `depth`.
fn at_depth(p: vec3<f32>, depth: f32) -> vec3<f32> {
// Looking down -z, so the plane is at -depth.
return p * (-depth / p.z);
}
@compute @workgroup_size(64, 1, 1)
fn cull_cs(@builtin(global_invocation_id) gid: vec3<u32>) {
let cells = clusters.grid.x * clusters.grid.y * clusters.grid.z;
let cell = gid.x;
if (cell >= cells) {
return;
}
let x = cell % clusters.grid.x;
let y = (cell / clusters.grid.x) % clusters.grid.y;
let z = cell / (clusters.grid.x * clusters.grid.y);
// The cell's four side rays, and the two depths that cap it.
let tile = vec2<f32>(
clusters.screen.x / f32(clusters.grid.x),
clusters.screen.y / f32(clusters.grid.y),
);
let min_ray = view_ray(vec2<f32>(f32(x), f32(y)) * tile);
let max_ray = view_ray(vec2<f32>(f32(x + 1u), f32(y + 1u)) * tile);
let near = clusters.screen.z;
let far = clusters.screen.w;
let ratio = far / near;
let near_depth = near * pow(ratio, f32(z) / f32(clusters.grid.z));
let far_depth = near * pow(ratio, f32(z + 1u) / f32(clusters.grid.z));
// The cell as a box around its eight corners. A froxel is not a box, so
// this is a little generous at the edges -- which costs a light in a
// cluster that did not quite need it, never a light missing from one that
// did.
let a = at_depth(min_ray, near_depth);
let b = at_depth(max_ray, near_depth);
let c = at_depth(min_ray, far_depth);
let d = at_depth(max_ray, far_depth);
let lo = min(min(a, b), min(c, d));
let hi = max(max(a, b), max(c, d));
var count = 0u;
let base = cell * MAX_LIGHTS_PER_CLUSTER;
for (var i = 0u; i < clusters.grid.w; i = i + 1u) {
if (count == MAX_LIGHTS_PER_CLUSTER) {
break;
}
let light = lights[i];
let centre = (view * vec4<f32>(light.position_range.xyz, 1.0)).xyz;
let range = light.position_range.w;
// Sphere against box: the distance from the centre to the nearest
// point of the box.
let nearest = clamp(centre, lo, hi);
let offset = centre - nearest;
if (dot(offset, offset) <= range * range) {
cluster_lights[base + count] = i;
count = count + 1u;
}
}
cluster_counts[cell] = count;
}