Edge Shadows Explained

The fringe

Key idea. Before PR #387, every shadowed silhouette carried an outline one pixel wide, darker or lighter than both surfaces it separated, because each edge pixel's farther surface was shaded with the nearer surface's shadow.

Ticket #353's visual audit captured every Reference scene at 2560×1440 on the M4 Pro and looked at them frame by frame. Along silhouettes it found pixels that were not a blend of the two surfaces they covered. Ticket #357 measured three, in 8-bit sRGB after tonemapping:

From ticket #357. The ticket puts a correct blend of the second row at about (65, 70, 77). After PR #387 the shadows edge pixel reads (158, 140, 134), between its neighbours.

The first row's edge pixel is brighter, in every channel, than the figure and the shadow on either side of it. The third is darker overall than both its neighbours. The second is as dark as the figure, darker in green and blue, where a blend with the lit ground would sit near (65, 70, 77). A blend of two colours lies between them, so these pixels were not blends: some of their samples carried the wrong shadow.

None of this was a surprise. When the shadow mask replaced the lit pass's own rays, Shadow Rays Explained wrote down what an edge pixel gives up: "a fringe one pixel wide at most, as strong as those fragments' share of the pixel's samples", and showed it in a live demo. This guide is the fix for that fringe.

One word per pixel

Key idea. MSAA stores colour per sample but shades once per pixel for each triangle, so anything the fragment shader reads per pixel has to know which surface it describes. The old shadow mask described only the nearest sample's.

With 4x MSAA, a pixel has four samples. The rasteriser decides per sample which triangle covers it, but runs the fragment shader once per pixel for each triangle that covers any of its samples, and writes that one colour into the samples the triangle covers. The resolve then averages the four.

That is the economy of MSAA (MSAA Fragment Merging Explained): a pixel inside a triangle shades once, and a pixel on an edge shades once per triangle, with each sample keeping its own triangle's colour. Edges come out antialiased because each surface's colour lands only in its own samples.

The shadow mask broke that bargain. Its compute pass ran one invocation per pixel, picked the pixel's nearest sample, the one with the greatest depth under reverse-Z (ADR 0015), traced every light's ray from the surface that sample shows, and wrote one word:

// shadow_mask.wgsl at 14e3446e, build_shadow_mask: one surface per pixel.
// The nearest sample has the greatest depth (ADR 0015); the first such.
let depth = max(max(depths.x, depths.y), max(depths.z, depths.w));

// main.wgsl at 14e3446e, fs_main: the same word for every fragment.
mask = textureLoad(shadow_mask, vec2<u32>(in.clip_position.xy), 0).rg;

Inside a surface, nothing is lost. At a silhouette, the pixel's samples show two surfaces, and the lit pass runs a fragment for each. Both read the word traced from the nearest surface. The farther surface's fragment then lights its samples with the nearer surface's verdict: where the sphere is lit and the ground behind it lies in shadow, the ground's samples there come out lit, and the other way round. The resolve averages right samples with wrong ones, and the pixel lands outside the range of its neighbours, as strongly as the farther surface's share of its samples.

The expensive cure

The textbook answer is per-sample shading: run the lit pass's fragment shader once per sample, so each sample can read a value of its own. Ticket #357 ruled it out on cost: ADR 0019's lit pass shades per pixel, and a forward pass shading per sample runs its fragment shader per sample at every pixel, edge or not, to fix about one pixel in seven on town. The fix had to keep one fragment per surface per pixel, and give each one the right word.

Surfaces, not triangles

Key idea. A pixel crossed by several triangles usually shows only one or two surfaces, so a small, fixed number of shadow words per pixel is enough. Praxis stores two.

The insight is Salvi and Vidimče's, from their 2012 paper on surface-based antialiasing: "even when a pixel is covered by multiple primitives it is generally only covered by a very small number of distinct surfaces." Two triangles of one mesh that meet inside a pixel continue the same surface, and their shading barely differs there. So, they write, "we can further reduce the per-pixel shading rate by computing only one color sample per surface", and "store only a small and fixed (e.g. 2 or 3) number of surfaces per pixel".

A shadow follows the same logic. Two triangles of a wall that meet inside a pixel see the same lights, and one ray from either answers for both. Only a discontinuity, such as a figure's outline over the ground behind it, puts two surfaces in one pixel that can answer differently. So the mask became a two-layer array, a word per surface:

// targets.rs at 7442717a.
/// The surfaces of a pixel the shadow mask holds a word for, a layer each:
/// the one its nearest sample shows, and the farthest other one where its
/// samples show more than one, as surface-based antialiasing stores a
/// small, fixed number of surfaces per pixel and shades each once
/// [ref:salvi-2012-sbaa-surfaces].
pub const SHADOW_MASK_SURFACES: u32 = 2;

• Layer 0 holds the nearest sample's surface, as the single layer did before. Every pixel writes it.
• Layer 1 holds a second surface's word, written only at pixels that have one. Each word is still an Rg32Uint pair: 59 light bits and, in the second half's top bits, the slice of the Froxel whose lights the bits follow (Froxel Lights Explained).
• shadow_mask_samples, an R32Uint target, holds four bits per pixel: which samples show the second surface.

All three live in RenderTargets with the list of chapter 6, so they are rebuilt on resize together with the bind groups that read them.

Finding the second surface

Key idea. A sample shows another surface when it lies off the nearest sample's triangle plane by more than one pixel's width at that distance. Measuring from the plane rather than comparing raw depths keeps a slanted floor one surface.

Deferred renderers with MSAA met this question first: which pixels need more than one shading sample? Lauritzen's 2010 SIGGRAPH course slides give the recipe: "Store G-buffer at sample frequency […] Only apply per-sample shading where necessary […] Want to find surface discontinuities […] Compare sample depths to depth derivatives". A depth derivative says how depth changes across the screen on the surface; extrapolated to another sample, it says what depth that sample would have if it showed the same surface. Praxis does the same with the nearest sample's triangle plane, which is what the derivative extrapolates:

// shadow_mask.wgsl at 7442717a, build_shadow_mask.
if depth > 0.0 {
    let point = vec2<f32>(pixel) + sample_offset(sample);
    let position = position_at(camera.view_proj, point, depth, target_size);
    let pixel_width = distance(
        position,
        position_at(camera.view_proj, point + vec2<f32>(1.0, 0.0), depth, target_size)
    );
    let plane_normal = octahedral_decode(textureLoad(surface_normals, pixel, sample).xy * 2.0 - 1.0);
    let off_plane = vec4<f32>(
        dot(position_at(camera.view_proj, vec2<f32>(pixel) + SAMPLE_OFFSETS[0], depths.x, target_size) - position, plane_normal),
        // … and samples 1 to 3 the same way
    );
    let second = abs(off_plane) > vec4<f32>(pixel_width) & depths > vec4<f32>(0.0);
    let bits = select(vec4<u32>(0u), vec4<u32>(1u, 2u, 4u, 8u), second);
    second_samples = bits.x | bits.y | bits.z | bits.w;

Why a plane, and why the triangle's

A floor seen at a grazing angle changes depth quickly from one sample to the next, though every sample shows the same floor. A test on raw depth differences would call those pixels edges and trace them twice for nothing; the demo in chapter 5 has that test as a switch, and on its sloping ground it lists 91 of its 112 pixels. Measuring each sample's distance from the nearest sample's plane cancels the slope: a sample on the same flat surface lies on the plane, whatever its depth.

The plane is the triangle's own. The depth passes write two normals per sample into the Surface normals (Shadow Terminator Explained): the triangle's, and the one interpolated over it for shading. The depths lie on the triangle, so the test takes the first pair, .xy. The interpolated normal bends to smooth the shading and would tilt the plane away from where the samples really are.

Why one pixel's width

A fixed distance in metres would suit one view distance only: a millimetre is huge for a pixel near the camera and nothing for a pixel on the horizon. The threshold is instead the world distance between the nearest sample and the point one pixel to its right at the same depth, the width of a pixel where the surface is. A sample off the plane by less than that is the same surface bending, or a neighbouring triangle of it; a sample farther off is something else.

The background, and which sample speaks for the second surface

A sample at depth 0, the far plane under reverse-Z, shows the background, which nothing shadows and whose word has no bits, so it is never a second surface. Where several samples are off the plane, the pass traces from the farthest of them, the least depth. It makes the other samples' depths 2, past the near plane's 1, so they never win the minimum:

// shadow_mask.wgsl at 7442717a.
// Depth 2 lies past the near plane's 1, so no other is less.
let candidates = select(vec4<f32>(2.0), depths, second);
let second_depth = min(min(candidates.x, candidates.y), min(candidates.z, candidates.w));
let second_sample = first_sample_of(candidates, second_depth);

Choosing the farthest is a rule for the rare pixel with three surfaces or more (chapter 8). With two, every off-plane sample shows the same surface, and any of them would do.

Which word a fragment reads

Key idea. Each fragment knows which of the pixel's samples it covers. If any of them is marked as the second surface's, it reads layer 1; otherwise layer 0.

A fragment shader can read its coverage: WGSL's sample_mask built-in input holds a bit for each sample of the pixel that the fragment's triangle covers. The pass wrote the second surface's samples as the same kind of bits, so one & decides:

// main.wgsl at 7442717a, fs_main.
fn fs_main(in: VertexOutput, @builtin(sample_mask) covered_samples: u32) -> @location(0) vec4<f32> {
    // …
    var mask = vec2<u32>(0u);
    if SHADOW_RAYS {
        let pixel = vec2<u32>(in.clip_position.xy);
        let second_samples = textureLoad(shadow_mask_samples, pixel, 0).r;
        let surface = select(0u, 1u, (covered_samples & second_samples) != 0u);
        mask = textureLoad(shadow_mask, pixel, surface, 0).rg;
    }

The nearer surface's fragment covers the nearest sample and its neighbours on the same plane, none of them marked, so it reads layer 0. The farther surface's fragment covers the marked samples, so it reads layer 1. Each colours its own samples with its own surface's shadow, and the resolve blends two right answers.

The word also carries its Froxel slice, the one at the sample it was traced from. The lit pass compares it with the slice at its own fragment and, where they differ, looks each light up in that Froxel's list. After the fix that happens, in the shader's words, only "where a slice ends between the sample and the pixel's centre, or at an edge of more than two surfaces".

The test at the seam

an_edge_pixel_resolves_between_a_shadowed_surface_and_a_lit_one, in src/render/renderer.rs, renders a panel whose edge crosses the middle of one pixel, so that the pixel's two left samples show the panel and its two right ones the floor behind. A square behind the camera casts the sun's shadow on the panel only, then on the floor only. In both cases the edge pixel's red channel must lie between its two neighbours', more than 16 steps from each, far past the tonemapper's dither. With the lit pass forced back to layer 0, the old behaviour, it fails: "the edge pixel 38 is not between the panel 38 and the floor 158".

Tracing a minority

Key idea. Second surfaces lie thinly along edges, so tracing them where they lie makes whole SIMD groups wait for a few lanes. Packing them into one frame-wide list, traced by an indirect dispatch the GPU sizes itself, costs roughly the lanes that need it.

On town, 13.7% of pixels have a second surface. The obvious code traces it in build_shadow_mask, behind an if. That cost town 4.1 ms.

Why a branch costs the whole group

A GPU runs invocations in lockstep groups that share one instruction stream: SIMD groups. Dougall Johnson's reverse-engineered reference to the G13, the M1's GPU, puts it plainly: "The G13 architecture has 32 threads per SIMD-group", and inactive threads are tracked in "a 32-bit execution mask". Praxis already sizes one workgroup for that width; culling.wgsl calls 32 invocations "a SIMD-group on Apple GPUs". When some lanes of a group take a branch, the whole group steps through it, the others masked off and waiting. Edges are thin lines, so they cross most groups: a group of 32 pixels with even one edge pixel among them runs a full second trace.

Lauritzen's notes describe the same trap for per-sample shading: "Lack of spatial locality causes hardware scheduling inefficiency […] Can use branching, but incurs scheduling problems […] Instead, reschedule per-sample pixels". Their tile-based recipe shades sample 0 of the whole tile, then packs "a list of pixels that require per-sample shading" and redistributes threads over it. Packing within each 8×8 or 16×16 workgroup did not help Praxis: PR #387 records 10.85 and 12.9 ms against 10.70. So the list is frame-wide.

The list

Each edge pixel claims an entry with one atomic add and writes its pixel and sample into it, packed into a u32: the column in the low 16 bits, the row in the next 14, the sample in the top 2. That is why RenderTargets::new asserts the target is at most 2¹⁴, 16,384, pixels tall. One atomic per pixel beat one per workgroup with barriers by 0.1 ms, and is simpler.

// shadow_mask.wgsl at 7442717a, build_shadow_mask.
let entry = atomicAdd(&second_surfaces.count, 1u);
second_surfaces.entries[entry] = invocation.x | (invocation.y << 16u) | (second_sample << 30u);
// The entry that opens a workgroup's run of the second dispatch
// adds that workgroup, in rows of MAX_WORKGROUPS_PER_DIMENSION.
let group_size = SHADOW_MASK_TILE_SIZE * SHADOW_MASK_TILE_SIZE;
if entry % group_size == 0u {
    let groups = entry / group_size + 1u;
    atomicMax(&second_surface_dispatch[0], min(groups, MAX_WORKGROUPS_PER_DIMENSION));
    atomicMax(
        &second_surface_dispatch[1],
        (groups + MAX_WORKGROUPS_PER_DIMENSION - 1u) / MAX_WORKGROUPS_PER_DIMENSION
    );
    atomicStore(&second_surface_dispatch[2], 1u);
}

• The GPU sizes the second dispatch. The CPU cannot know how many pixels will be listed, so the second dispatch is indirect: it reads its workgroup counts from a buffer the first dispatch writes. The entry that opens each run of 64 adds a workgroup with atomicMax, so the counts end at the number of workgroups the list needs.
• Rows of 65,535. A dispatch may have at most 65,535 workgroups along each dimension, MAX_WORKGROUPS_PER_DIMENSION. At 64 entries a workgroup, a list past 65,535 × 64, about 4.2 million pixels, would overflow x, and a target of more pixels than that, such as 3840×2160's 8.3 million, can list that many. A review fix folds the count into rows: x is at most 65,535, y the rows needed, and each invocation works out its entry from both.
• One read of the count per workgroup. Every invocation of trace_second_surfaces must check its entry against the count. Lane 0 loads it and workgroupUniformLoad shares it, because, as the shader says, "an atomic every invocation loads from one address serialises them".
• One pass, cleared first. Both dispatches run in one compute pass. Before it, the encoder clears the dispatch buffer and the list's count, its first 4 bytes.
• A stand-in buffer. wgpu checks the buffers a dispatch uses as one set, and an indirect dispatch may not read its arguments from a buffer that a bound group can write. The first dispatch's group binds the dispatch buffer read-write, so the second binds a 12-byte stand-in in that slot, same layout; one function builds both groups.

// shadow_mask.wgsl at 7442717a.
@compute @workgroup_size(SHADOW_MASK_TILE_SIZE * SHADOW_MASK_TILE_SIZE)
fn trace_second_surfaces(
    @builtin(workgroup_id) group: vec3<u32>,
    @builtin(num_workgroups) groups: vec3<u32>,
    @builtin(local_invocation_index) lane: u32
) {
    if lane == 0u {
        listed_count = atomicLoad(&second_surfaces.count);
    }
    let entry = (group.y * groups.x + group.x) * SHADOW_MASK_TILE_SIZE * SHADOW_MASK_TILE_SIZE + lane;
    if entry >= workgroupUniformLoad(&listed_count) {
        return;
    }
    // … decode the pixel and sample, load its depth, trace into layer 1
}

The result: tracing the list cost town 1.08 ms of GPU frame, against 4.1 ms in place.

What it costs

Key idea. Judge the fix by the GPU frame, 1.08 ms more on town, not by the lit pass's timestamp, which rose 7.9 ms only because Metal now runs it alongside the shadow mask pass's second dispatch.

town on the M4 Pro at 2560×1440 against main's Baselines, from praxis report, in µs:

From PR #387's Metrics, rows trimmed as the PR trims them, its p95 and Budget columns left out. town's GPU frame Budget is 12,000 µs, which its Baseline already exceeded.

Why the lit pass did not get 7.9 ms slower

On Metal, passes can overlap, so a pass's own timestamps can include time spent waiting on the pass before (Frame Timing Explained). Metal now overlaps the lit pass's start with the shadow mask pass's indirect dispatch, so the lit pass's span takes in that wait. A run without the second dispatch read the lit pass at 5.85 ms. So neither pass's own delta is the fix's cost. Only spans count shared time once: the GPU frame rose 1.08 ms, and the shadow-ray time, the rays-on frame minus the rays-off frame, rose by nearly the same, 1.09 ms. ADR 0019's amendment quotes 1.15 ms from another run; about 1.1 ms either way.

The gate, which ADR 0019 times with shadow rays off, moved 19 µs, within noise: with rays off, the shadow mask pass does not run.

Memory, allocations, aliasing

• 16 bytes per pixel. Layer 1, 8 bytes; shadow_mask_samples, 4; a list entry, 4, since every pixel could be listed. With the list's 4-byte count, that is 58,982,404 bytes at 2560×1440, as ADR 0019 records. The bench's gpu_memory_bytes rose by 58,982,428, the two 12-byte dispatch buffers included.
• Heap allocations: the bench's heap_allocations rose by 9, from wgpu's validation of the indirect dispatch.
• Aliasing RMSE −13%. The bench compares the Golden's frame with the same frame rendered at three times the size and box-filtered in linear light. aliasing_rmse fell from 6,413 to 5,579, a ratio of 0.870, and aliasing_p99 from 20,967 to 18,210. The fringe was aliasing, and the counter saw it go.

Limits

Key idea. Two words per pixel cover a silhouette over one background. Three surfaces in a pixel, coverage taken before the depth test, and the jagged edges of shadows themselves remain.

• Three surfaces or more. Where a pixel's samples show three surfaces, every off-plane sample is marked and the farthest is traced, so the middle surface's fragment takes the farthest one's word. ADR 0019 records it. SBAA suggests storing 2 or 3 surfaces per pixel; Praxis stores 2.
• Coverage before the depth test. sample_mask may include samples where the fragment then fails the depth test. A fragment of the nearest surface whose triangle runs under a nearer surface at another sample then counts that sample as covered, and reads the second word. PR #387 deferred it: revive it if a visual audit finds a fringe at steep interpenetrating geometry, or if wgpu exposes coverage after the depth test.
• The shadow's own edge. The mask is still one ray per light per surface per pixel, so the boundary between shadow and light across a surface aliases at one sample per pixel while geometry gets 4x MSAA. That is ticket #358, waiting on a literature study, #364. Keeping the mask per surface means #358 can read it per surface.

Check yourself

A floor seen at a grazing angle has samples whose depths differ by more than a pixel's width. Why is it not listed for a second trace?

The test does not compare depths. It measures each sample's distance from the plane of the nearest sample's triangle, along that triangle's normal. Every sample of a flat floor lies on that plane, at distance zero, whatever its depth. Only a sample off the plane by more than one pixel's world width at the nearest sample's distance counts as another surface.

At an edge pixel, shadow_mask_samples holds 0b1100. A fragment's sample_mask is 0b1000. Which word does it read, and which did it read before PR #387?

0b1000 & 0b1100 is not zero, so it covers a second-surface sample and reads layer 1, traced from the farthest of samples 2 and 3. Before, every fragment read the one word, traced from the nearest sample, which is on the other surface: the fragment shaded its samples with the other surface's shadow.

Only 13.7% of town's pixels need a second trace. Why did tracing them in place cost 4.1 ms, and why does the lit pass's +7.9 ms not mean the fix costs 7.9 ms?

Edge pixels lie along thin lines that cross most 32-lane SIMD groups, and a group with one lane in the branch runs the whole trace, so in place most groups traced twice. A frame-wide list packs them into as few groups as they fill: 1.08 ms of GPU frame. The lit pass's timestamp grew because Metal overlaps its start with the second dispatch, so its span includes that work; a run without the dispatch read it at 5.85 ms. The GPU frame, a span, counts the overlap once.

Reading list

Grouped by source, each with what it is good for here. Code links are pinned to 7442717a.

Glossary

Shadow mask, Surface normals and Froxel are CONTEXT.md's definitions, the first two shortened; so are the tooltips for Golden, Baseline and Reference scene. The rest are this guide's own.