Metal #14 Post-processing

This note covers post-processing: the chain of screen-space operations applied to a rendered image, and how to wire it together in Metal.

HDR and tone mapping

Post-processing presupposes rendering into an HDR intermediate, not straight into an 8-bit drawable.

Swiftlet d = MTLTextureDescriptor.texture2DDescriptor(
    pixelFormat: .rgba16Float, width: w, height: h, mipmapped: false)
d.usage = [.renderTarget, .shaderRead, .shaderWrite]

The reason is simple: real-world luminance ranges far beyond [0, 1]. The sun is orders of magnitude brighter than white paper. Clip everything at 1 during rendering and bloom becomes impossible (every overexposed region is the same 1.0) and tone mapping becomes meaningless.

Tone mapping compresses the HDR range into what a display can show. The simplest is Reinhard:

MSLfloat3 reinhard(float3 c) { return c / (1.0 + c); }

Mathematically clean, but highlights come out washed out, because it compresses every luminance uniformly and lacks the soft shoulder that film has. The current default is an ACES curve fit:

MSLfloat3 ACESFilm(float3 x) {
    const float a = 2.51, b = 0.03, c = 2.43, d = 0.59, e = 0.14;
    return saturate((x * (a * x + b)) / (x * (c * x + d) + e));
}

It preserves shadow contrast, rolls off smoothly in the highlights, and includes a slight hue shift that mimics film. This is close to the industry default now.

Multiply by an exposure value before tone mapping. Auto-exposure can be built by mipping a luminance image and reading the last level for the average, with a temporal low-pass filter so the image does not pulse.

Bloom

Bloom simulates the halo real lenses produce around bright sources. Three steps: extract, blur, add back. That flow is covered in detail in Unity Shader #13 Post-Processing; here I will stick to what is Metal-specific.

The current approach is not a single-radius Gaussian but progressive downsampling followed by progressive upsampling:

Plain Text1/2 ──> 1/4 ──> 1/8 ──> 1/16
                          │
1/2 <── 1/4 <── 1/8 <──────┘   (each level added to the same-size downsample result)

Each level uses a small 13-tap kernel, once down and once up. The result has very wide reach — equivalent to a huge blur radius — for the cost of a few small-kernel samples, and because halos at several scales are summed, it looks considerably more natural than one Gaussian.

Watch out for fireflies when downsampling: an isolated very bright pixel becomes a visible square block. The fix is a luminance-weighted average (the Karis average) in the first downsample:

MSLfloat weight(float3 c) { return 1.0 / (1.0 + luminance(c)); }
// weighted average rather than arithmetic mean

SSAO

Screen-space ambient occlusion estimates, from the depth buffer, how much each pixel is occluded by nearby geometry.

The basic algorithm: take several samples in the normal-oriented hemisphere around a pixel, project them back to screen space, and compare each sample's depth against the depth buffer. If the recorded depth is nearer, that direction is occluded.

MSLkernel void ssao(texture2d<float> depthTex [[texture(0)]],
                 texture2d<float> normalTex [[texture(1)]],
                 texture2d<float, access::write> output [[texture(2)]],
                 constant SSAOParams &params [[buffer(0)]],
                 uint2 gid [[thread_position_in_grid]])
{
    float3 posVS = reconstructViewPos(gid, depthTex, params);
    float3 n     = normalTex.read(gid).xyz * 2.0 - 1.0;

    float occlusion = 0.0;
    for (uint i = 0; i < params.sampleCount; i++) {
        float3 samplePos = posVS + orientedHemisphereSample(i, n, gid) * params.radius;
        float2 sampleUV  = projectToUV(samplePos, params.projection);
        float  sceneZ    = linearDepth(depthTex.sample(s, sampleUV).r, params);

        // range check: distant geometry must not occlude a nearby pixel
        float rangeCheck = smoothstep(0.0, 1.0, params.radius / abs(posVS.z - sceneZ));
        occlusion += (sceneZ >= samplePos.z + params.bias ? 1.0 : 0.0) * rangeCheck;
    }
    output.write(1.0 - occlusion / params.sampleCount, gid);
}

That rangeCheck is essential. Without it, the edge of a near object casts a false dark halo onto a distant wall, because samples separated by a large depth gap still count as occlusion.

The sample count can never be large, so SSAO output is inherently noisy. The standard treatment is a 4×4 random rotation texture to push the noise to high frequency, followed by a 4×4 blur to remove it. Sixteen samples then produce something close to the quality of 256.

Compute or a fullscreen triangle?

Post-processing traditionally uses a fullscreen triangle and a fragment shader. But a compute shader is usually better here:

Threadgroup memory is available. In a blur kernel, the samples needed by neighbouring threads overlap heavily. Cooperatively loading a block into threadgroup memory and reading from there cuts texture bandwidth several times over:

MSLkernel void blur_h(texture2d<float> src [[texture(0)]],
                   texture2d<float, access::write> dst [[texture(1)]],
                   uint2 gid [[thread_position_in_grid]],
                   uint2 lid [[thread_position_in_threadgroup]])
{
    threadgroup float4 cache[64 + 2 * RADIUS];

    // cooperative load, including the halo on both sides
    cache[lid.x + RADIUS] = src.read(gid);
    if (lid.x < RADIUS) {
        cache[lid.x] = src.read(uint2(gid.x - RADIUS, gid.y));
        cache[lid.x + 64 + RADIUS] = src.read(uint2(gid.x + 64, gid.y));
    }
    threadgroup_barrier(mem_flags::mem_threadgroup);

    float4 sum = 0;
    for (int i = -RADIUS; i <= RADIUS; i++)
        sum += cache[lid.x + RADIUS + i] * kernelWeights[i + RADIUS];
    dst.write(sum, gid);
}

Several outputs at once. One kernel can write multiple textures without the render-target count and format-agreement constraints.

Real scatter. A fragment shader can only write its own pixel; a compute shader can write anywhere, which is what makes histograms and luminance reductions possible.

No render pass load/store cost. A compute encoder has no attachments, so there is no tile load and no writeback.

The cost is losing fixed-function blending and rasterizer optimizations. For simple gather-only effects — tone mapping, colour grading — a fullscreen triangle may still be faster.

Organizing the chain

A typical chain:

Plain TextHDR scene ─> SSAO ─> bloom (down/up) ─> motion blur ─> tonemap + grading ─> FXAA/TAA ─> drawable

Manage the intermediates with the MTLHeap described in Metal #9 Render Passes so non-overlapping results share memory. And fold tone mapping and colour grading into one pass, expressing the whole colour transform as a 3D LUT: that saves one fullscreen read and write, which at 4K is not a small amount.

The next note covers reflection and refraction.