21. Metal Performance Shaders

Written by Marius Horga

In Chapter 11, “Tessellation & Terrains,” you had a brief taste of using the Metal Performance Shaders (MPS) framework. MPS consists of low-level, fine-tuned, high-performance kernels that run off the shelf with minimal configuration. In this chapter, you’ll dive a bit deeper into the world of MPS.

Overview

The MPS kernels make use of data-parallel primitives that are written in such a way that they can take advantage of each GPU family’s characteristics. The developer doesn’t have to care about which GPU the code needs to run on, because the MPS kernels have multiple versions of the same kernel written for every GPU you might use. Think of MPS kernels as convenient black boxes that work efficiently and seamlessly with your command buffer. Simply give it the desired effect, a source and destination resource (buffer or texture), and then encode GPU commands on the fly!

The Sobel filter is a great way to detect edges in an image. In the projects folder for this chapter, open and run sobel.playground and you’ll see such an effect (left: original image, right: Sobel filter applied):

Assuming you already created a device object, a command queue, a command buffer and a texture object for the input image, there are only two more lines of code you need to apply the Sobel filter to your input image:

let shader = MPSImageSobel(device: device)
shader.encode(commandBuffer: commandBuffer, 
              sourceTexture: inputImage,
              destinationTexture: drawable.texture)

MPS kernels are not thread-safe, so it’s not recommended to run the same kernel on multiple threads that are all writing to the same command buffer concurrently.

Moreover, you should always allocate your kernel to only one device, because the kernel’s init(device:) method could allocate resources that are held by the current device and might not be available to another device.

Note: MPS kernels provide a copy(with:device:) method that allows them to be copied to another device.

The MPS framework serves a variety of purposes beyond image filters. One of those areas is Neural Networks, which is not covered in this book. Instead, you’ll stay focused on:

1. Image processing
2. Matrix/vector mathematics
3. Ray tracing

Image processing

There are a few dozen MPS image filters, among the most common being:

• Morphological (area min, area max, dilate, erode).
• Convolution (median, box, tent, Gaussian blur, Sobel, Laplacian, and so on).
• Histogram (histogram, histogram equalization, histogram specification).
• Threshold (binary, binary inverse, to zero, to zero inverse, and so on).
• Manipulation (conversion, Lanczos scale, bilinear scale, transpose).

Note: For a complete list of MPS kernels, consult Apple’s official page at https://developer.apple.com/documentation/metalperformanceshaders/image_filters. If you want to implement your own filters, you can get inspired from Gimp’s list at https://docs.gimp.org/en/filters.html.

An RGB image is nothing but a matrix with numbers between 0 and 255 (when using 8-bit color channels). A greyscale image only has one such matrix because it only has one channel. For color images, there are three separate RGB channels (red, green, blue) so consequently three matrices, one for each channel.

One of the most important operations in image processing is convolution which is an operation consisting of applying a much smaller matrix, often called the kernel, to the original image and obtaining the desired effect as a result.

As an example, this matrix is used for obtaining Gaussian blur:

Note: You can find a list of common kernels at https://en.wikipedia.org/wiki/Kernel_(image_processing)

Here’s a diagram showing how the kernel is applied to two pixels:

And here’s how the result shown in green was calculated:

(6 * 1  +  7 * 2  +  3 * 1  + 
 4 * 2  +  9 * 4  +  8 * 2  + 
 9 * 1  +  2 * 2  +  3 * 1) / 16 = 6

In this example, 16 represents the weight, and it’s not a randomly chosen number — it’s the sum of the numbers from the convolution kernel matrix.

When you need to apply convolution to the border pixels, you can apply padding to the input matrix. For example, when the center of a 3×3 convolution kernel overlaps with the image element at position (0, 0), the image matrix needs to be padded with an extra row and extra column of zeros. However, if the bottom rightmost element of the convolution kernel overlaps with the image element at position (0, 0), the image matrix needs to be padded with two extra rows and two extra columns of zeros.

Applying the convolution kernel to an image matrix padded with an extra row, and extra column of zeros, gives you this calculation for a 3×3 kernel:

(0 * 1  +  0 * 2  +  0 * 1  + 
 0 * 2  +  6 * 4  +  7 * 2  + 
 0 * 1  +  4 * 2  +  9 * 1) / 9 = 6

In this case, the weight, 9, is the sum of the numbers from the convolution kernel matrix that are affecting only the non-zero numbers from the image matrix (4 + 2 + 2 + 1).

Something like this is straightforward, but when you need to work with larger kernels and multiple images that need convolution, this task might become non-trivial.

You know how to calculate by hand and apply convolution to an image — and at the very beginning of the chapter you saw an MPS filter on an image too — but how about using MPS in your engine?

What if you were to implement bloom in your engine?

Guess what? You are going to do just that next!

Bloom

The bloom effect is quite a spectacular one. It amplifies the brightness of objects in the scene and makes them look luminous as if they’re emitting light themselves.

Below is a diagram that gives you an overview of how to achieve bloom:

Here are the steps you’re going to take:

• Render the entire scene to a texture.
• Apply a threshold filter to this texture. This will amplify the lighter parts of the image, making them brighter.
• Apply a blur filter to the threshold texture from the previous step.
• Combine this texture with the initial scene for the final image.

The project

In the starter folder, open the Bloom project. Build and run, and you’ll see a familiar scene from previous chapters.

The project currently renders to the view’s drawable texture. Instead of sending this texture straight to the screen, you’ll intercept and use the drawable texture as input to the threshold filter.

First, import the MPS framework at the top of Renderer.swift:

import MetalPerformanceShaders

Define two textures at the top of Renderer:

var outputTexture: MTLTexture!
var finalTexture: MTLTexture!

outputTexture will hold the blurred threshold texture, and finalTexture will hold this texture combined with the initial render.

Add this to the end of mtkView(_:drawableSizeWillChange:). Remember, you call this method at the end of init(metalView), and every time the window resizes:

outputTexture = 
    createTexture(pixelFormat: view.colorPixelFormat,
                  size: size)
finalTexture = 
    createTexture(pixelFormat: view.colorPixelFormat,
                  size: size)

You create the MTLTextures using a helper method. You’ll be able to read, write and render to these textures.

Image Threshold To Zero

Using the following test, MPSImageThresholdToZero is a filter that returns either the original value for each pixel having a value greater than a specified brightness threshold or 0:

destinationColor = sourceColor > thresholdValue 
                     ? sourceColor : 0

This filter has the effect of making darker areas black, while the lighter areas retain their original color value.

In draw(in:), look for // MPS brightness filter, and below that line, add the following:

let brightness =
    MPSImageThresholdToZero(device: Renderer.device,
                            thresholdValue: 0.2,
                            linearGrayColorTransform: nil)
brightness.label = "MPS brightness"
brightness.encode(commandBuffer: commandBuffer,
                  sourceTexture: drawable.texture,
                  destinationTexture: outputTexture)

Here, you create an MPS kernel to create a threshold texture with a custom brightness threshold set to 0.2 — where all pixels with less than a color value of 0.2 will be turned to black. The input texture is the drawable texture, which contains the current rendered scene. The result of the filter will go into outputTexture.

Internally, the MPS kernel samples from drawable.texture, so you have to set the drawable to be used for read/write operations. Add this to the top of init(metalView:):

metalView.framebufferOnly = false

Metal optimizes drawable as much as possible, so setting framebufferOnly to false will affect performance slightly.

To be able to see the result of this filter, you’ll blit outputTexture back into drawable.texture. You should be familiar with the blit encoder from Chapter 14, “Multipass & Deferred Rendering.”

Locate // blit encoder in draw(in:), and add this afterward:

finalTexture = outputTexture
guard let blitEncoder = commandBuffer.makeBlitCommandEncoder() 
      else { return }
let origin = MTLOriginMake(0, 0, 0)
let size = MTLSizeMake(drawable.texture.width, 
                       drawable.texture.height, 
					   1)
blitEncoder.copy(from: finalTexture, sourceSlice: 0, 
                 sourceLevel: 0,
                 sourceOrigin: origin, sourceSize: size,
                 to: drawable.texture, destinationSlice: 0,
                 destinationLevel: 0, destinationOrigin: origin)
blitEncoder.endEncoding()

This copies the output of the previous filter into the drawable texture.

Build and run the project, and you’ll see the brightness texture.

Notice how only half of the tree and part of the train were bright enough to make it to this texture. These white areas are all you need to create the bloom effect.

Before using this texture, you need to add a little fuzziness to it which will make the model edges appear to glow. You can accomplish this with another MPS kernel — the Gaussian blur.

Gaussian blur

MPSImageGaussianBlur is a filter that convolves an image with a Gaussian blur with a given sigma value (the amount of blur) in both the X and Y directions.

In Renderer.swift, in draw(in:), locate MPS blur filter, and add this code:

let blur = MPSImageGaussianBlur(device: Renderer.device,
                                sigma: 9.0)
blur.label = "MPS blur"
blur.encode(commandBuffer: commandBuffer,
            inPlaceTexture: &outputTexture,
            fallbackCopyAllocator: nil)

In-place encoding is a special type of encoding where, behind the curtains, the input texture is processed, stored to a temporary texture and finally written back to the input texture without the need for you to designate an output texture.

The fallbackCopyAllocator argument allows you to provide a closure where you can specify what will happen to the input image should the in-place normal encoding fail.

Build and run the project, and you’ll see the result of this blur.

Image add

The final part of creating the bloom effect is to add the pixels of this blurred image to the pixels of the original render.

MPSImageArithmetic, as its name suggests, performs arithmetic on image pixels. Subclasses of this include MPSImageAdd, MPSImageSubtract, MPSImageMultiply and MPSImageDivide.

Adding the rendered scene pixels to the lighter blurred pixels will brighten up those parts of the scene. In contrast, adding them to the black pixels will leave them unchanged.

Just before // blit encoder, add this:

let add = MPSImageAdd(device: Renderer.device)
add.encode(commandBuffer: commandBuffer, 
           primaryTexture: drawable.texture, 
           secondaryTexture: outputTexture, 
           destinationTexture: finalTexture)

This adds the drawable texture to outputTexture and places the result in finalTexture.

Since you’re now creating finalTexture by combining two other textures, remove the line:

finalTexture = outputTexture

Build and run the project, and you’ll see this:

Notice how the tree, train tank and wheel caps appear to glow more? Awesome bloom!

Matrix/vector mathematics

You learned in the previous section how you could quickly apply a series of MPS filters that are provided by the framework. But what if you wanted to make your own filters?

You can create your own filter functions and calculate convolutions yourself, however, when working with large matrices and vectors, the amount of math involved might get overwhelming.

The MPS framework not only provides image processing capability, but it also provides functionality for decomposition and factorizing matrices, solving systems of equation and multiplying matrices and/or vectors on the GPU in a fast, highly parallelized fashion. You’re going to look at matrix multiplication next.

Create a new empty playground for macOS, named matrix.playground, and add the following code to it:

import MetalPerformanceShaders

guard let device = MTLCreateSystemDefaultDevice(),
      let commandQueue = device.makeCommandQueue() 
else { fatalError() }

let size = 4
let count = size * size

guard let commandBuffer = commandQueue.makeCommandBuffer() 
else { fatalError() }

commandBuffer.commit()
commandBuffer.waitUntilCompleted()

This creates a Metal device, command queue, command buffer and adds a couple of constants you’ll need later.

Above the line where you create the command buffer, add a new method that lets you create MPS matrices:

func createMPSMatrix(withRepeatingValue: Float) -> MPSMatrix {
  // 1
  let rowBytes = MPSMatrixDescriptor.rowBytes(
                                   forColumns: size,
                                   dataType: .float32)
  // 2
  let array = [Float](repeating: withRepeatingValue, 
                      count: count)
  // 3
  guard let buffer = device.makeBuffer(bytes: array,
                                       length: size * rowBytes,
                                       options: []) 
  else { fatalError() }
  // 4
  let matrixDescriptor = MPSMatrixDescriptor(
                                   rows: size,
                                   columns: size,
                                   rowBytes: rowBytes,
                                   dataType: .float32)
                                             
  return MPSMatrix(buffer: buffer, descriptor: matrixDescriptor)
}

Going through the code:

1. Retrieve the optimal number of bytes between one row and the next. Whereas simd matrices expect column-major order, MPSMatrix uses row-major order.
2. Create a new array and populate it with the value provided as an argument to this method.
3. Create a new buffer with the data from this array.
4. Create a matrix descriptor; then create the MPS matrix using this descriptor and return it.

Use this new method to create and populate three matrices. You’ll multiply A and B together and place the result in C. Add this just before creating the command buffer:

let A = createMPSMatrix(withRepeatingValue: 3)
let B = createMPSMatrix(withRepeatingValue: 2)
let C = createMPSMatrix(withRepeatingValue: 1)

Add this to create a MPS matrix multiplication kernel:

let multiplicationKernel = MPSMatrixMultiplication(
                              device: device,
                              transposeLeft: false,
                              transposeRight: false,
                              resultRows: size,
                              resultColumns: size,
                              interiorColumns: size,
                              alpha: 1.0,
                              beta: 0.0)

Below the line where you create the command buffer, add this code to encode the kernel:

multiplicationKernel.encode(commandBuffer:commandBuffer,
                            leftMatrix: A,
                            rightMatrix: B,
                            resultMatrix: C)

You multiply A and B together and the result is placed in C. At the very end of the playground add this code to read C:

// 1
let contents = C.data.contents()
let pointer = contents.bindMemory(to: Float.self, 
                                  capacity: count)
// 2
(0..<count).map {
  pointer.advanced(by: $0).pointee
}

Going through the code:

1. Read the result back from the matrix C into a buffer typed to Float, and set a pointer to read through the buffer.

2. Create an array filled with the values from the buffer.

Run the playground; click Show Result on the last line; right-click on the graph of the results, and then choose “Value History.”

You’ll see that the array contains 16 values, all of which are the number 24.0. That’s because the matrix is of size 4×4, and multiplying one row of A with one column of B results in the value 24.0, which is 2×3 added four times.

This is only a small matrix, but you can change the size of the matrix in the size variable at the top of the playground, and the matrix multiplication will still be blisteringly fast.

Ray tracing

In Chapter 18, “Rendering with Rays,” you looked briefly at ray tracing and path tracing. In this section of the chapter, you’re going to implement an MPS-accelerated raytracer, which is, in fact, a path tracer variant using the Monte Carlo integration.

Note: You’ll create a scene using an adaptation of Apple’s sample app https://developer.apple.com/documentation/metalperformanceshaders/metal_for_accelerating_ray_tracing which is translated to Swift, modified to work with .obj files for scene objects and simplified for better understanding.

Remember that with the Monte Carlo integration you shoot primary rays for each pixel, and when there’s a hit in the scene, you shoot one more secondary ray in a random direction for each primary ray shot.

Ray tracing and path tracing are known for photorealistic rendering and accurate shadows, reflection, refraction, ambient occlusion, area lights, depth of field, and so on.

The path tracing algorithm looks like this:

For each pixel on the screen:
  Reset the pixel color C.
    For each sample (random direction): 
      Shoot a ray and trace its path.
      C += incoming radiance from ray.
    C /= number of samples 

The outline of this section is as follows:

1. Primary rays
2. Shadow rays
3. Secondary rays

All right, that’s a plan. Ready, set, go!

1. Primary rays

Primary rays render the equivalent of a rasterized scene, but the most expensive part of ray tracing is finding all of the intersections between rays and triangles in the scene.

The MPS framework provides a high performance MPSRayIntersector class specifically created to accelerate ray-triangle intersection tests on the GPU.

The MPSRayIntersector object uses two inputs: a ray buffer and an acceleration structure. It outputs into another buffer all the intersections it finds for each ray cast.

MPSTriangleAccelerationStructure is the class used to build the acceleration structure from vertices that describe the triangles in a scene. You pass the structure on to the intersector.

1.0 The starter app

Time for some coding! Open the starter project named Raytracing, build and run it. Although you won’t see anything but a dull solid color, the starter project contains much of the setup needed.

In RendererExtension.swift, loadAsset(name:position:scale) is the method that loads object positions, normals and colors into separate arrays. Call this method with OBJ files to add objects to the scene.

In Renderer:

• createScene() loads a default scene.
• createBuffers() creates buffers from the OBJ file arrays, sets up the uniforms buffer, and creates a random buffer that will contain random numbers.
• update() gets called every frame. It updates the uniforms and then generates 256 random numbers between 0 and 1. You’ll use these random numbers for antialiasing, choosing a random point on the light source and bouncing secondary rays randomly.
• draw(in:) has sections that you’ll fill out for the ray tracing, and it renders a simple quad at the end of the method. It’s this quad that’s currently colored turquoise in the fragment shader.

1.1 Create the render target

As you go through the various passes, you’ll write to a render target texture. You’ll accumulate values, and this will be the texture you render onto the screen quad.

In Renderer.swift, at the top of Renderer, add the render target property:

var renderTarget: MTLTexture!

In mtkView(_:drawableSizeWillChange:) create the texture by adding this to the end of the method:

let renderTargetDescriptor = MTLTextureDescriptor()
renderTargetDescriptor.pixelFormat = .rgba32Float
renderTargetDescriptor.textureType = .type2D
renderTargetDescriptor.width = Int(size.width)
renderTargetDescriptor.height = Int(size.height)
renderTargetDescriptor.storageMode = .private
renderTargetDescriptor.usage = [.shaderRead, .shaderWrite]
renderTarget = device.makeTexture(descriptor: renderTargetDescriptor)

This sets up the texture with a descriptor that defines it for both read and write operations on the GPU.

1.2 Create the Ray Intersector

When you generate the primary rays in a kernel, you send the results to a Ray struct array of a particular format. The ray intersector decides this format.

In Renderer.swift, at the top of the class, declare the ray intersector object:

var intersector: MPSRayIntersector!
let rayStride = 
MemoryLayout<MPSRayOriginMinDistanceDirectionMaxDistance>.stride 
  + MemoryLayout<float3>.stride

rayStride specifies how large the Ray struct will be. It also allows for holding custom fields in the struct. As well as holding origin, minimum distance, direction and maximum distance, you’ll also hold a custom float3 color field.

Add a new method to Renderer to create the intersector:

func buildIntersector() {
  intersector = MPSRayIntersector(device: device)
  intersector?.rayDataType 
      = .originMinDistanceDirectionMaxDistance
  intersector?.rayStride = rayStride
}

rayDataType matches the stride you just set up and determines what fields the ray buffer structure should contain. Add a call to this method at the end of init(metalView:):

buildIntersector()

1.3 Generate primary rays

Before you can generate the primary rays, you need to create a new compute pipeline state and a buffer to hold the generated rays. At the top of Renderer, add this code:

var rayPipeline: MTLComputePipelineState!
var rayBuffer: MTLBuffer!
var shadowRayBuffer: MTLBuffer!

At the same time as you create and manage the ray buffer, you’ll also set up the shadow ray architecture — which is similar to the ray architecture — but you’ll handle the shadow calculations later.

Add this to the end of mtkView(_:drawableSizeWillChange:):

let rayCount = Int(size.width * size.height)
rayBuffer = device.makeBuffer(length: rayStride * rayCount,
                              options: .storageModePrivate)
shadowRayBuffer = 
    device.makeBuffer(length: rayStride * rayCount,
                      options: .storageModePrivate)

This creates the ray and shadow buffers large enough to accommodate one ray for each pixel of the rendered image.

In buildPipelines(view:), add this code before the do statement:

let computeDescriptor = MTLComputePipelineDescriptor()
computeDescriptor.threadGroupSizeIsMultipleOfThreadExecutionWidth 
    = true

You set threadGroupSizeIsMultipleOfThreadExecutionWidth to true to tell the compiler to optimize the compute kernel. For that to work, you need to also set the thread group size to be a multiple of threadExecutionWidth when you dispatch threads to do work.

Inside the do statement, add this code:

computeDescriptor.computeFunction = library.makeFunction(
                                           name: "primaryRays")
rayPipeline = try device.makeComputePipelineState(
                                 descriptor: computeDescriptor,
                                 options: [],
                                 reflection: nil)

Here, you set up the pipeline for the generate primary rays compute kernel.

In draw(in:), add this code right below // MARK: generate rays:

// 1
let width = Int(size.width)
let height = Int(size.height)
let threadsPerGroup = MTLSizeMake(8, 8, 1)
let threadGroups = 
    MTLSizeMake((width + threadsPerGroup.width - 1)
                                  / threadsPerGroup.width,
                (height + threadsPerGroup.height - 1)
                                  / threadsPerGroup.height,
                 1)
// 2
var computeEncoder = commandBuffer.makeComputeCommandEncoder()
computeEncoder?.label = "Generate Rays"
computeEncoder?.setBuffer(uniformBuffer, 
                          offset: uniformBufferOffset,
                          index: 0)
computeEncoder?.setBuffer(rayBuffer, offset: 0, index: 1)
computeEncoder?.setBuffer(randomBuffer, 
                          offset: randomBufferOffset,
                          index: 2)
computeEncoder?.setTexture(renderTarget, index: 0)
computeEncoder?.setComputePipelineState(rayPipeline)
computeEncoder?.dispatchThreadgroups(threadGroups,
  threadsPerThreadgroup: threadsPerGroup)
computeEncoder?.endEncoding()

Going through the code:

1. Configure the number of threads per group and the number of thread groups that will run on the GPU.
2. Create a compute encoder for the ray pipeline, send the texture and buffers to the GPU and dispatch the threads to execute the kernel function on the GPU.

To generate primary rays, you launch a 2D grid of threads, one per render target pixel. Each thread will write a Ray struct to the ray buffer in the kernel. The intersector will read this array of Rays.

Open Raytracing.metal. It contains a few helper functions that you’ll need later. At the top of the file, after // add structs here, add this:

struct Ray {
  packed_float3 origin;
  float minDistance;
  packed_float3 direction;
  float maxDistance;
  float3 color;
};

The intersector’s rayDataType specifies that the struct should be of type .originMinDistanceDirectionMaxDistance, so you define the struct according to that type. You also define the extra custom field color which will later hold the scene objects’ color.

Each primary ray starts at the camera position (origin) and passes through a pixel on the image plane resulting in one primary ray per pixel.

Add this kernel function below Ray:

kernel void 
    primaryRays(constant Uniforms & uniforms [[buffer(0)]],
             device Ray *rays [[buffer(1)]],
             device float2 *random [[buffer(2)]],
             texture2d<float, access::write> t [[texture(0)]],
             uint2 tid [[thread_position_in_grid]]) {
  // 1
  if (tid.x < uniforms.width && tid.y < uniforms.height) {
    // 2
    float2 pixel = (float2)tid;
    float2 r = random[(tid.y % 16) * 16 + (tid.x % 16)];
    pixel += r;
    float2 uv = 
        (float2)pixel / float2(uniforms.width, uniforms.height);
    uv = uv * 2.0 - 1.0;
    // 3
    constant Camera & camera = uniforms.camera;
    unsigned int rayIdx = tid.y * uniforms.width + tid.x;
    device Ray & ray = rays[rayIdx];
    ray.origin = camera.position;
    ray.direction = 
        normalize(uv.x * camera.right + uv.y * camera.up 
                     + camera.forward);
    ray.minDistance = 0;
    ray.maxDistance = INFINITY;
    ray.color = float3(1.0);
    // 4
    t.write(float4(0.0), tid);
  }
}

Going through the code:

1. Check to make sure you’re not using threads outside of the pixel grid.

2. Assign one thread to each pixel on the UV plane. Randomize the pixel slightly to prevent aliasing.

3. Create a Ray for each thread, and fill out the Ray struct. Set the ray’s origin to the camera position, calculate the ray direction, set the min/max distance the ray could extend to, and finally, give the ray a white color.

4. Reset the render target texture to a black image. You’ll add color later.

At this point, you could use a refreshment!

1.4 Accumulation

You’re writing to the render target texture that you’ll combine with the other textures; you’ll create these other textures later for shadows and secondary rays, and then render to the background quad. You’ll set this render up now so that you can see your progress.

First, set up the pipeline state and final render target texture. In Renderer.swift, add this code at the top of Renderer:

var accumulatePipeline: MTLComputePipelineState!
var accumulationTarget: MTLTexture!

In buildPipelines(view:), inside the do statement create the pipeline state:

computeDescriptor.computeFunction = library.makeFunction(
  name: "accumulateKernel")
accumulatePipeline = try device.makeComputePipelineState(
  descriptor: computeDescriptor, options: [], reflection: nil)

In mtkView(_:drawableSizeWillChange:), add this code at the end of the method to create the texture:

accumulationTarget = device.makeTexture(
  descriptor: renderTargetDescriptor)

In draw(in:), locate // MARK: accumulation, and add this code right below it:

computeEncoder = commandBuffer.makeComputeCommandEncoder()
computeEncoder?.label = "Accumulation"
computeEncoder?.setBuffer(uniformBuffer, 
                          offset: uniformBufferOffset,
                          index: 0)
computeEncoder?.setTexture(renderTarget, index: 0)
computeEncoder?.setTexture(accumulationTarget, index: 1)
computeEncoder?.setComputePipelineState(accumulatePipeline)
computeEncoder?.dispatchThreadgroups(threadGroups,
  threadsPerThreadgroup: threadsPerGroup)
computeEncoder?.endEncoding()

Here, you call the accumulation kernel which writes from renderTarget to accumulationTarget.

A little further down in draw(in:), add this line before the draw call:

renderEncoder.setFragmentTexture(accumulationTarget, index: 0)

This sends the texture to the fragment shader for the final draw.

In Shaders.metal, add this kernel function:

kernel void accumulateKernel(constant Uniforms & uniforms,
                   texture2d<float> renderTex,
                   texture2d<float, access::read_write> t,
                   uint2 tid [[thread_position_in_grid]])
{
  if (tid.x < uniforms.width && tid.y < uniforms.height) {
    // 1
    float3 color = renderTex.read(tid).xyz;
    if (uniforms.frameIndex > 0) {
      // 2
      float3 prevColor = t.read(tid).xyz;
      prevColor *= uniforms.frameIndex;
      color += prevColor;
      color /= (uniforms.frameIndex + 1);
    }
    t.write(float4(color, 1.0), tid);
  }
}

Here, you’re outputting an average of the input texture over the previous frames.

Going through the code:

1. Read the color value at each pixel from the current render target texture, and if the frame index is 0, simply return the same color and store it into the accumulation texture.
2. If the frame index is not 0, add the current color to the color value from the current render target texture and store it into the accumulation texture.

In fragmentShader, replace the return line with this code:

constexpr sampler s(min_filter::nearest,
                    mag_filter::nearest,
                    mip_filter::none);
float3 color = tex.sample(s, in.uv).xyz;
return float4(color, 1.0);

Instead of returning turquoise, you sample the color from the accumulation texture and return that color as the current pixel color.

Build and run the app, and you’ll get a black screen. Remember, you set the render target texture to black during the ray generation kernel.

You’ve now generated random rays and set up the interceptor. The other part of interception is the acceleration structure created from the model vertices.

1.5 Create the acceleration structure

In Renderer.swift, at the top of the class, declare the acceleration structure object:

var accelerationStructure: MPSTriangleAccelerationStructure!

Add a new method to create the structure:

func buildAccelerationStructure() {
  accelerationStructure = 
    MPSTriangleAccelerationStructure(device: device)
  accelerationStructure?.vertexBuffer = vertexPositionBuffer
  accelerationStructure?.triangleCount = vertices.count / 3
  accelerationStructure?.rebuild()
}

This creates the acceleration structure from the provided vertex buffer. loadAsset(name:position:scale:) loads up all of the vertices for the models into the vertex position buffer, so the number of triangles in the acceleration structure is the number of vertices divided by 3.

Add a call to this method to the end of init(metalView:):

buildAccelerationStructure()

1.6 Intersect Rays with the Scene

The next stage is to take the generated rays, and the acceleration structure, and use the intersector to combine them into an intersection buffer that contains all of the hits where a ray coincides with a triangle.

First, you need a new buffer to store the intersections. In Renderer.swift, add this to the top of Renderer:

var intersectionBuffer: MTLBuffer!
let intersectionStride = 
MemoryLayout<MPSIntersectionDistancePrimitiveIndexCoordinates>.stride

Similar to setting up the Ray struct that contains the generated rays, you’ll define an Intersection struct to hold the generated intersections. intersectionStride defines the stride of this struct.

In mtkView(_:drawableSizeWillChange:), add this code at the end to set up the buffer:

intersectionBuffer = device.makeBuffer(
  length: intersectionStride * rayCount,
  options: .storageModePrivate)

In draw(in:), add this code right below // MARK: generate intersections between rays and model triangles:

intersector?.intersectionDataType = .distancePrimitiveIndexCoordinates
intersector?.encodeIntersection(
  commandBuffer: commandBuffer,
  intersectionType: .nearest,
  rayBuffer: rayBuffer,
  rayBufferOffset: 0,
  intersectionBuffer: intersectionBuffer,
  intersectionBufferOffset: 0,
  rayCount: width * height,
  accelerationStructure: accelerationStructure)

You specify the data type to match the stride. Your intersection struct for the kernel will contain distance, primitive indices and coordinates.

The intersector’s encodeIntersection method computes intersections and encodes its results to a Metal command buffer.

For primary rays, you set the intersection type to MPSIntersectionType.nearest so that the intersector returns the intersections that are closest to the camera. Then, you send the intersector the generated rays, the acceleration structure and the intersection buffer to receive the intersection results.

1.7 Use intersections for shading

The last step in casting primary rays is shading. This depends on intersection points and vertex attributes, so yet another compute kernel applies the lighting based on this information. At the top of Renderer, add a new pipeline for this new kernel:

var shadePipelineState: MTLComputePipelineState!

In buildPipelines(view:), inside the do statement, create the pipeline state:

computeDescriptor.computeFunction = library.makeFunction(
  name: "shadeKernel")
shadePipelineState = try device.makeComputePipelineState(
  descriptor: computeDescriptor,options: [], reflection: nil)

In draw(in:), under // MARK: shading, add this code for the shading compute encoder:

computeEncoder = commandBuffer.makeComputeCommandEncoder()
computeEncoder?.label = "Shading"
computeEncoder?.setBuffer(uniformBuffer, 
                          offset: uniformBufferOffset,
                          index: 0)
computeEncoder?.setBuffer(rayBuffer, offset: 0, index: 1)
computeEncoder?.setBuffer(shadowRayBuffer, offset: 0, index: 2)
computeEncoder?.setBuffer(intersectionBuffer, offset: 0, 
                          index: 3)
computeEncoder?.setBuffer(vertexColorBuffer, offset: 0, 
                          index: 4)
computeEncoder?.setBuffer(vertexNormalBuffer, offset: 0, 
                          index: 5)
computeEncoder?.setBuffer(randomBuffer, 
                          offset: randomBufferOffset,
                          index: 6)
computeEncoder?.setTexture(renderTarget, index: 0)
computeEncoder?.setComputePipelineState(shadePipelineState!)
computeEncoder?.dispatchThreadgroups(threadGroups,
  threadsPerThreadgroup: threadsPerGroup)
computeEncoder?.endEncoding()

Compared to the first compute encoder, you’re now also sending to the GPU: the shadow ray buffer, the intersection buffer, the vertex color buffer and the vertex normal buffer. You also switched to using the shadePipeline state.

In Raytracing.metal, add this struct after Ray:

struct Intersection {
  float distance;
  int primitiveIndex;
  float2 coordinates;
};

The contents of this struct depend upon the intersectionDataType that you specified in the intersector.

As you can see, intersections depend on the distance between the intersecting ray origin and the geometry, the primitive index, and the barycentric coordinates of the intersection on the triangle.

The shading kernel runs one thread per pixel and has the same effect as a fragment shader; however, you’ll have to do the interpolation of vertex attributes yourself.

In Raytracing.metal, uncomment the function interpolateVertexAttribute. This helper function performs the color and normal interpolations that the rasterizer would normally do for you.

template<typename T>
inline T interpolateVertexAttribute(device T *attributes, 
                     Intersection intersection) {
  // 1
  float3 uvw;
  uvw.xy = intersection.coordinates;
  uvw.z = 1.0 - uvw.x - uvw.y;
  // 2
  unsigned int triangleIndex = intersection.primitiveIndex;
  T T0 = attributes[triangleIndex * 3 + 0];
  T T1 = attributes[triangleIndex * 3 + 1];
  T T2 = attributes[triangleIndex * 3 + 2];
  return uvw.x * T0 + uvw.y * T1 + uvw.z * T2;
}

Going through the code:

1. Get the first two barycentric coordinates from the intersection buffer and compute the third one.
2. Get the vertex attributes from the vertex attribute buffer using primitive index offsets, and then return the interpolated vertex attribute.

The shading kernel itself is pretty hefty, but you’ll create it gradually. Add the new kernel at the end of Raytracing.metal:

kernel void shadeKernel(uint2 tid [[thread_position_in_grid]],
                        constant Uniforms & uniforms,
                        device Ray *rays,
                        device Ray *shadowRays,
                        device Intersection *intersections,
                        device float3 *vertexColors,
                        device float3 *vertexNormals,
                        device float2 *random,
                        texture2d<float, access::write> renderTarget)
{
  if (tid.x < uniforms.width && tid.y < uniforms.height) {
	
  }
}

• rays holds a ray color for every sample.
• You’ll use shadowRays shortly to calculate the shadow.
• intersections holds the array of intersections in the form of the intersection data type struct Intersection.

From these, you’ll be able to work out the intersection point for the thread, so you also pass vertex normals and vertex colors. You’ll interpolate the intersection result to get the correct normal and color for the thread.

• The texture renderTarget is unnecessary for the generation of rays, but you’ll temporarily write the result of this shading kernel to it so you can see what’s happening.

Add this inside the if statement:

unsigned int rayIdx = tid.y * uniforms.width + tid.x;
device Ray & ray = rays[rayIdx];
device Ray & shadowRay = shadowRays[rayIdx];
device Intersection & intersection = intersections[rayIdx];
float3 color = ray.color;

This extracts the primary ray, shadow ray and intersection per thread/pixel and gets the primary ray color. On the first run, this will be white, which you initially set in primaryRays.

Continue adding after the previous code:

// 1
if (ray.maxDistance >= 0.0 && intersection.distance >= 0.0) {
  float3 intersectionPoint = ray.origin + ray.direction
                              * intersection.distance;
  float3 surfaceNormal = 
      interpolateVertexAttribute(vertexNormals,
                                 intersection);
  surfaceNormal = normalize(surfaceNormal);
  // 2
  float2 r = random[(tid.y % 16) * 16 + (tid.x % 16)];
  float3 lightDirection;
  float3 lightColor;
  float lightDistance;
  sampleAreaLight(uniforms.light, r, intersectionPoint,
                  lightDirection, lightColor, lightDistance);
  // 3                
  lightColor *= saturate(dot(surfaceNormal, lightDirection));
  color *= interpolateVertexAttribute(vertexColors, 
                                      intersection);
}
else {
  ray.maxDistance = -1.0;
}
// 4
renderTarget.write(float4(color, 1.0), tid);

Going through the code:

1. If both the ray maximum distance and the intersection distance are non-negative, calculate the intersection point and the surface normal using the interpolateVertexAttribute function with the vertex normal buffer.
2. Use another utility function named sampleAreaLight that takes in a light object, a random direction and an intersection point, and returns the light direction, color and distance.
3. Adjust the light color using the surface normal and light direction from the previous steps. Also, adjust the pixel color by using the interpolateVertexAttribute function again, but this time, with the vertex color buffer.
4. Write the calculated color for the pixel to the render target texture.

Build and run the project, and you’ll see the results of your hard work so far. Instead of using the rasterizer to render your objects, you created the scene with rays and testing vertex intersections.

So far you’re not calculating shadows, but it’s a start.

Important! Remove this from the end of shadeKernel:

renderTarget.write(float4(color, 1.0), tid);

That line was just for visualization, and you shouldn’t write to the render target until you’ve also calculated the shadow rays.

2. Shadow rays

As well as calculating the color of the pixel in the final texture, you’ll need to check if the point is in shadow.

You’ll cast a ray from the point to the light source and check whether the shadow ray reaches the light. If it doesn’t, then the point is in the shadow.

In the previous kernel, you sent a buffer named shadowRays which you’re currently not using. You’ll write the shadow result into this buffer.

In Raytracing.metal, in shadeKernel, locate color *= interpolateVertexAttribute(vertexColors, intersection);. Add this afterwards:

shadowRay.origin = intersectionPoint + surfaceNormal * 1e-3;
shadowRay.direction = lightDirection;
shadowRay.maxDistance = lightDistance - 1e-3;
shadowRay.color = lightColor * color;

This saves the origin of the point, the direction from the point to the light and the current color. You’ll use these values in the next kernel.

In the else part of the same conditional, reset the maximum distance when both the ray maximum distance and the intersection distance are non-negative:

shadowRay.maxDistance = -1.0;

The following diagram illustrates the similarity to the primary rays step. You cast shadow rays from the intersection points to the light source.

If the ray does not reach the light source, it will be in the shadow:

In Renderer.swift, add a new pipeline state for the new kernel. Add this code at the top of Renderer:

var shadowPipeline: MTLComputePipelineState!

In buildPipelines(view:), inside the do statement add this code:

computeDescriptor.computeFunction = library.makeFunction(
                                           name: "shadowKernel")
shadowPipeline = 
    try device.makeComputePipelineState(
                      descriptor: computeDescriptor,
                      options: [],
                      reflection: nil)

This creates the pipeline for the upcoming shadow kernel.

Shadow rays need a maximum distance so they don’t go beyond the light source origin. Primary rays don’t need that, but they do need the triangle index or barycentric coordinates, neither of which are needed anymore here. Shadow rays carry the color from the shading kernel to the final kernel.

You can reuse the intersector and acceleration structure that you used for primary rays to compute shadow ray intersections, but you’ll need to configure it to use a ray data type that supports shadow rays based on the differences listed earlier.

In draw(in:), locate // MARK: shadows, and add this code right below:

intersector?.label = "Shadows Intersector"
intersector?.intersectionDataType = .distance
intersector?.encodeIntersection(
                commandBuffer: commandBuffer,
                intersectionType: .any,
                rayBuffer: shadowRayBuffer,
                rayBufferOffset: 0,
                intersectionBuffer: intersectionBuffer!,
                intersectionBufferOffset: 0,
                rayCount: width * height,
                accelerationStructure: accelerationStructure!)

Note that you’re now using shadowRayBuffer for the intersection generation. Because shadows don’t require the triangle index and coordinates anymore, you set the intersection data type to MPSIntersectionDataType.distance. You can still use the same Intersection struct in the kernel, but the other fields will be ignored by the MPSRayIntersector.

For primary ray intersections you had to know the nearest surface to the camera that intersects the ray, but now it doesn’t matter which surface intersects a shadow ray. If any triangles exist between the primary intersection and the light source, the primary intersection is shadowed so when computing shadow ray intersections set the intersector’s intersection type to .any.

Under the previous code, add the compute encoder for the shadow kernel:

computeEncoder = commandBuffer.makeComputeCommandEncoder()
computeEncoder?.label = "Shadows"
computeEncoder?.setBuffer(uniformBuffer, 
                          offset: uniformBufferOffset,
                          index: 0)
computeEncoder?.setBuffer(shadowRayBuffer, offset: 0, index: 1)
computeEncoder?.setBuffer(intersectionBuffer, offset: 0, 
                          index: 2)
computeEncoder?.setTexture(renderTarget, index: 0)
computeEncoder?.setComputePipelineState(shadowPipeline!)
computeEncoder?.dispatchThreadgroups(
                   threadGroups,
                   threadsPerThreadgroup: threadsPerGroup)
computeEncoder?.endEncoding()

Add this new kernel to the end of Raytracing.metal:

kernel void shadowKernel(uint2 tid [[thread_position_in_grid]],
             constant Uniforms & uniforms,
             device Ray *shadowRays,
             device float *intersections,
             texture2d<float, access::read_write> renderTarget)
{
  if (tid.x < uniforms.width && tid.y < uniforms.height) {
    // 1
    unsigned int rayIdx = tid.y * uniforms.width + tid.x;
    device Ray & shadowRay = shadowRays[rayIdx];
    float intersectionDistance = intersections[rayIdx];
    // 2
    if (shadowRay.maxDistance >= 0.0 
          && intersectionDistance < 0.0) {
      float3 color = shadowRay.color;
      color += renderTarget.read(tid).xyz;
      renderTarget.write(float4(color, 1.0), tid);
    }
  }
}

Going through the code:

1. Get the current thread index and create both a shadow ray and intersection distance at the current pixel.
2. If the shadow ray’s maximum distance is non-negative, but the distance to the intersection point is negative, add the shadow color to the current color from the render target texture, and save it back to the render target texture. If the shadow ray’s intersection distance is negative, it means the intersection point wasn’t in shadow because it reached the light source.

Build and run the project, and you’ll finally see a shaded scene.

3. Secondary rays

This scene looks quite dark because you’re not bouncing any light around. In the real world, light bounces off all surfaces in all directions.

To reproduce this effect, you’ll iterate over the central kernels several times. Each time, you’ll send the secondary rays in a random direction; this adds diffuse reflected light to areas in shadow.

In Renderer, in draw(in:), extend a for loop from // generate intersections between rays and model triangles to just above // accumulation:

for _ in 0..<3 {
    // MARK: generate intersections between rays and model triangles
    // MARK: shading
    // MARK: shadows
}  
// MARK: accumulation

This repeats the kernels three times. The more iterations you have, the clearer and more vibrant the image becomes over time, however, the final render will take much longer to complete.

On the first iteration, the ray direction is taken from the initial primaryRays kernel; but for secondary rays, the ray direction, or bounce, should be random.

In Raytracing.metal, update the secondary rays by giving them a random direction.

In shadeKernel, locate where you calculated shadowRay in the if statement, and add the following code afterward:

float3 sampleDirection = sampleCosineWeightedHemisphere(r);
sampleDirection = alignHemisphereWithNormal(sampleDirection,
                                            surfaceNormal);
ray.origin = intersectionPoint + surfaceNormal * 1e-3f;
ray.direction = sampleDirection;
ray.color = color;

The project already has the functions sampleCosineWeightedHemisphere and alignHemisphereWithNormal added. They’re responsible for the random direction of the secondary rays and for reducing the amount of noise from the rendered image.

Build and run the project. You should see this image:

Where to go from here?

What a great journey this has been. In this chapter, you were able to use the MPS framework to:

• Add Bloom effect to the scene.
• Write a matrix multiplication basic playground.
• Create an entirely ray-traced scene.

Apple documentation for the MPS frameworks is still being written, but one particular section to pay attention to is the Tuning Hints at https://developer.apple.com/documentation/metalperformanceshaders/tuning_hints.

Also, for more information about using the MPS framework for image processing, matrix multiplication and raytracing, you can read references.markdown included with this chapter.