27. GPU Command Encoding

Written by Marius Horga & Caroline Begbie

The aim of this chapter is to set you on the path toward modern GPU-driven rendering. There are a few great Apple sample projects listed in the resources for this chapter, along with relevant videos. However, the samples can be quite intimidating. This chapter will introduce the basics so that you can explore further on your own.

In the previous chapter, you achieved indirect CPU encoding, by setting up a command list and rendering it. You created a loop that executes serially on the CPU. This loop is one that you can easily parallelize.

Each ICB draw call command executes one after another, but by moving the command creation loop to the GPU, you can create each command at the same time over multiple GPU cores:

When you come to write real-world apps, setting up the render loop at the very start of the app is impractical. In each frame, you’ll be determining which models to render. Are the models in front of the camera? Is the model occluded by another model? Should you render a model with lower level of detail? By creating the command list every frame, you have complete flexibility in which models you should render, and which you should ignore.

As you’ll see, the GPU is amazingly fast at creating these render command lists, so you can include this process each frame.

The Starter Project

➤ In Xcode, open the starter project, and build and run the app.

The starter project is almost the same as the final project from the previous chapter with these exceptions:

• The radio button options are both for indirect encoding, one on the GPU and one on the CPU.
• The two render passes are held in IndirectRenderPass.swift and GPURenderPass.swift. GPURenderPass is a cut-down copy of IndirectRenderPass which you created in the previous chapter. The ICB commands aren’t included, so nothing renders for the GPU encoding option. You’ll add the commands in a shader function that runs on the GPU.
• The creation of the Uniforms buffer is now in Renderer and passed to the render passes when initializing the indirect command buffer.

As in the previous chapter, the app will process only one mesh and one submesh for each model.

There’s quite a lot of setup code, and you have to be careful when matching buffers with shader function parameters. If you make an error, it’s difficult to debug it, and your computer may lock up. Running the app on an external device, such as iPhone or iPad is preferable, if slightly slower.

These are the steps you’ll take through this chapter:

1. Organize your scene data.
2. Add the scene data to one big buffer.
3. Create the compute shader function.
4. Create the compute pipeline state object.
5. Encode the ICB.
6. Set up the compute shader threads and arguments.

1. Organizing Your Scene

Instead of handing the GPU one model at a time to encode, you’ll give a GPU compute shader function your whole scene organized into buffers. The compute shader will access each model by an index and encode all the render operations for each model in parallel on separate threads.

For the models’ vertex buffers and materials, you’ll create argument buffers containing the GPU addresses for all the necessary model buffers. You’ll need to hold :

• Mesh data: the addresses of the vertex buffers and submesh indices. For simplicity, your app only processes one submesh. In a real app, you’d separate this out into mesh data and submesh data.
• Model data: the materials and transform data for each model.

This will make up the scene data, so, just for fun, you’ll have an extra step of indirection by having a scene argument buffer that points to the mesh and model data.

➤ In the Shaders folder, create a new header file called SceneData.h.

In this file, you’ll define the necessary structures for both Metal and Swift.

➤ Before #endif, add this code:

#if __METAL_VERSION__
// MARK: - Metal Shading Language

#include <metal_stdlib>
using namespace metal;

struct SceneData {
  constant float3* positionsAndNormals;
  constant float2* uvs;
  constant uint32_t* indices;
  uint32_t indexType;
  uint32_t indexCount;
  constant ShaderMaterial* materials;
  constant ModelParams* modelParams;
};

#else 
// MARK: - Swift side

# endif

Here, you create the structures that you’ll need for accessing the argument buffers in Metal.

➤ After // Mark: - Swift side, add this code:

#include <Metal/Metal.h>

struct SceneData {
  uint64_t positions;
  uint64_t uvs;
  uint64_t indices;
  uint32_t indexType;
  uint32_t indexCount;
  uint64_t materials;
  uint64_t modelParams;
};

Here, you define the Swift resource argument buffers that point to the scene data. These match the Metal resources.

➤ Open Common.h, and after defining ModelParams, add this code:

#import "SceneData.h"

You need to add it after ModelParams, as SceneData requires knowledge of this structure.

2. Creating the Scene Data Buffer

Your current model data structure looks like this:

Ideally, you’d flatten your data into several buffers with this scene hierarchy:

However, as you’re creating a less complex app with only one mesh and one submesh per model, you’ll flatten all your data into one sceneBuffer that holds all the data for all the models:

➤ Open GPURenderPass.swift in the Render Passes folder and add these new properties to GPURenderPass that you’ll use to hold your scene data:

var sceneBuffer: MTLBuffer!
var modelParamsBufferArray: [MTLBuffer] = []

You’ll initialize these buffers in initialize(models:). Notice that initializeICBCommands(_:) differs from the previous chapter. It now only consists of setting up the indirect command buffer that the compute shader will fill.

➤ At the end of initialize(models:), add this code:

let sceneBufferSize = MemoryLayout<SceneData>.stride * models.count
sceneBuffer = Renderer.device.makeBuffer(length: sceneBufferSize)!
sceneBuffer.label = "Scene Buffer"
var scenePtr = sceneBuffer.contents()
  .assumingMemoryBound(to: SceneData.self)
for model in models {
  let mesh = model.meshes[0]
  let submesh = mesh.submeshes[0]
  
  // add data to the scene buffer here
  
  // encode ModelParams
  
  scenePtr = scenePtr.advanced(by: 1)
}

You initialize the scene buffer with the correct size. You then set up a pointer binding the memory to SceneData so you can access the contents more easily.

Iterating through all the models, you’ll now add the data to the scene buffer.

➤ Replace // add data to the scene buffer here with this code:

scenePtr.pointee.positions = mesh.vertexBuffers[0].gpuAddress
scenePtr.pointee.uvs = mesh.vertexBuffers[1].gpuAddress
scenePtr.pointee.indices = submesh.indexBuffer.gpuAddress
scenePtr.pointee.indexType = submesh.indexType == .uint16 ? 0 : 1
scenePtr.pointee.indexCount = UInt32(submesh.indexCount)
scenePtr.pointee.materials = model.meshes[0].submeshes[0]
  .materialBuffer.gpuAddress

Each buffer resource has a GPU address. Apple Silicon has a unified memory architecture, where both CPU and GPU share the same physical memory. The advantage is that there is little transfer overhead.

The model’s vertex buffers and submesh data are the details that the GPU will need to render each model.

Adding the model’s transform and tiling data is a little more complex. You’ll still be using the vertex function vertex_main and the fragment function fragment_main to process the render. These functions expect a structure ModelParams. However, the compute shader can’t create a new buffer from a structure. You’ll need to transfer ModelParams to a buffer, and then add this buffer to the scene buffer.

➤ Replace // encode ModelParams with this code:

// 1
var modelParams = ModelParams(
  modelMatrix: model.transform.modelMatrix,
  tiling: model.tiling)
// 2
let modelParamsBufferSize = MemoryLayout<ModelParams>.stride
let modelParamsBuffer = Renderer.device.makeBuffer(
  bytes: &modelParams, length: modelParamsBufferSize)!
modelParamsBuffer.label = "Model Params"
// 3
scenePtr.pointee.modelParams = modelParamsBuffer.gpuAddress
// 4
modelParamsBufferArray.append(modelParamsBuffer)

Going through the code:

1. Fill out ModelParams with the data that the vertex and fragment functions expect.
2. Create an MTLBuffer from ModelParams.
3. Add the GPU address pointer to the ModelParams buffer to the scene buffer.
4. Retain the buffer in memory by adding it to modelParamsBufferArray. If you don’t do this, the app will release modelParamsBuffer as soon as it has finished using it in the for loop.

You’ve now made a massive scene buffer that you can transfer to the GPU with one binding.

3. Creating the Compute Shader Function

Now you’ll create the indirect command buffer on the GPU. Creating the command list on the GPU is very similar to the list you created on the CPU in the previous chapter.

➤ In the Shaders folder, create a new Metal file named ICB.metal, and add the following:

#import "Common.h"

// 1
struct ICBContainer {
  command_buffer icb [[id(0)]];
};

kernel void encodeICB(
  // 2
  constant SceneData* scene [[buffer(0)]],
  constant Uniforms &uniforms [[buffer(UniformsBuffer)]],
  // 3
  device ICBContainer *icbContainer [[buffer(ICBBuffer)]],
  // 4
  uint modelIndex [[thread_position_in_grid]])
{
}

Going through the code:

1. You can only transfer an indirect command buffer to the GPU via an argument buffer. In GPURenderPass, you’ll encode icb into a container buffer shortly.
2. Here is the data for the scene. The compute function will extract each model’s data from scene and use it to fill out the draw command.
3. The indirect command buffer needs to be in the device space, as you’ll be writing to it in this function.
4. The compute function will process every model on its own thread, and the position in the grid will provide the index into scene.

Now that you’ve set up all the data, it’s an easy task to encode the draw call.

➤ Add this to encodeICB:

// 1
SceneData model = scene[modelIndex];
command_buffer icb = icbContainer->icb;

// 2
bool isVisible = true;
// 3
render_command cmd(icb, modelIndex);
if (isVisible) {
  cmd.set_vertex_buffer(&uniforms, UniformsBuffer);
  cmd.set_vertex_buffer(model.positionsAndNormals, VertexBuffer);
  cmd.set_vertex_buffer(model.uvs, UVBuffer);
  cmd.set_vertex_buffer(model.modelParams, ModelParamsBuffer);
  cmd.set_fragment_buffer(model.materials, MaterialBuffer);
  cmd.set_fragment_buffer(model.modelParams, ModelParamsBuffer);
} else {
// 4
  cmd.reset();
}

Going through the code:

1. You retrieve the model and draw arguments using the thread position in grid.
2. isVisible is doing a lot of heavy lifting here. You may have wondered what you’re gaining from moving the encoding to the GPU. This is the place where you can decide whether or not to render the model. You can call a function to work out whether the model is behind the camera. If the model has multiple levels of detail, you could work out which one to render.
3. As you’re not doing any visibility testing here, you always create the render command and encode the operations just as you did in Swift.
4. If you don’t want to render this particular model, you tell the ICB to ignore this draw.

Finally, you’ll encode the draw call.

➤ Add this code before the else in encodeICB:

if (model.indexType == 0) {
  // uint16 indices
  cmd.draw_indexed_primitives(
    primitive_type::triangle,
    model.indexCount,
    (constant ushort*) model.indices,
    1);
} else {
  // uint32 indices
  cmd.draw_indexed_primitives(
    primitive_type::triangle,
    model.indexCount,
    (constant uint32_t*) model.indices,
    1);
}

Here, you create the draw call, testing which index type the model is using. In your app, the ground model uses uint32 indices and the house model uint16. It’s very important to get this data type right, otherwise the vertex function won’t be able to access the indices correctly, and you’ll get weird visual errors that are hard to debug.

You’ve now encoded a complete draw call, and that’s all that’s required for the compute function. Your next task is to set up the compute function on the CPU side, with a compute pipeline state and pass all the data to the compute function.

4. Creating the Compute Pipeline State Object

➤ Open GPURenderPass.swift, and create these new properties in GPURenderPass:

let icbPipelineState: MTLComputePipelineState
let icbComputeFunction: MTLFunction

To run the compute function you just created, you’ll need a new compute pipeline state.

➤ Add the following code to the end of init():

icbComputeFunction =
  Renderer.library.makeFunction(name: "encodeICB")!
icbPipelineState = PipelineStates.createComputePSO(
  function: "encodeICB")

This code creates the compute function in the Metal library, and also the compute pipeline state.

5. Encoding the ICB

The encodeICB compute function requires as input a buffer that contains the indirect command buffer.

➤ In GPURenderPass, add the container buffer:

var icbContainer: MTLBuffer!

➤ At the end of initializeICBCommands(_:), add this code:

let icbEncoder = icbComputeFunction.makeArgumentEncoder(
  bufferIndex: ICBBuffer.index)
icbContainer = Renderer.device.makeBuffer(
  length: icbEncoder.encodedLength,
  options: [])
icbEncoder.setArgumentBuffer(icbContainer, offset: 0)
icbEncoder.setIndirectCommandBuffer(icb, index: 0)

You can’t send an indirect command buffer directly to the GPU, as it needs to be verified internally first as being suitable for the GPU. You create the argument encoder with reference to the compute function that will use it. So when you set the argument buffer as the container, together with the indirect command buffer, this verification can take place.

6. Setting up the Compute Command Encoder

You’ve done all the preamble and setup code. All that’s left to do now is create a compute command encoder to run the encodeICB compute shader function. The function will create a render command to render every model.

➤ Still in GPURenderPass.swift, add a new method to GPURenderPass:

func encodeICB(
  commandBuffer: MTLCommandBuffer,
  models: [Model],
  uniforms: MTLBuffer
) {
  guard let computeEncoder = 
    commandBuffer.makeComputeCommandEncoder() else { return }
  computeEncoder.label = "GPU Encoding"
  
  computeEncoder.setComputePipelineState(icbPipelineState)
  computeEncoder.setBuffer(sceneBuffer, offset: 0, index: 0)
  computeEncoder.setBuffer(
    uniforms, offset: 0, index: UniformsBuffer.index)
  computeEncoder.setBuffer(
    icbContainer, offset: 0, index: ICBBuffer.index)
}

Here, you create the compute command encoder and set the arguments that the compute function encodeICB will use.

➤ Add the following code to the end of encodeICB(commandBuffer:models:uniforms:):

// Dispatch threads
let threadExecutionWidth = icbPipelineState.threadExecutionWidth
let drawCount = models.count // should be number of draw calls
let threads = MTLSize(width: drawCount, height: 1, depth: 1)
let threadsPerThreadgroup = MTLSize(
  width: threadExecutionWidth, height: 1, depth: 1)
computeEncoder.dispatchThreads(
  threads, threadsPerThreadgroup: threadsPerThreadgroup)
computeEncoder.endEncoding()

Here, you decide on how many draw calls you need to create and dispatch the compute encoder with that many threads.

➤ Call this method at the top of draw(commandBuffer:scene:uniforms:):

encodeICB(
  commandBuffer: commandBuffer,
  models: scene.models,
  uniforms: uniforms)

Now you’re ready to run the app. You already set up the render command encoder in the previous chapter, and you can use the same execution command on the ICB. The only difference is that you filled the ICB on the GPU instead.

➤ Build and run the app, and your scene will appear.

Because all your buffers are implicitly bound to the GPU, the GPU capture is unable to put together the scene, and you’ll get a blank render target there.

➤ Back in GPURenderPass.swift, add this code to the end of useResources(encoder:models:), before encoder.popDebugGroup:

encoder.useResource(sceneBuffer, usage: .read, stages: [.vertex, .fragment])
modelParamsBufferArray.forEach {
  encoder.useResource($0, usage: .read, stages: [.vertex, .fragment])
}

Note: even though useResource(_:,usage:,stages:) is very efficient, and adds little to the run time load, you can work out which buffers are actually resident, and surround the others with #if DEBUG...endif. That way you’ll be able to view the GPU capture, and the commands won’t appear in the release version.

➤ Now, build and run the app, and capture the GPU frame. Take a look at the Command Buffer, the GPU Encoding Pass, and select dispatchThreads to see all the resources bound to the compute pass.

➤ Double click Scene Buffer under Compute and check that your scene data is properly formatted.

You should have two entries, one for each model. Each of the properties has an arrow indicating that it points to another buffer. You can click this arrow to view the contents of the other buffers.

The GPU Command Encoding render pass resources should be the same as the previous chapter.

As in the previous chapter, the app probably doesn’t show much speed improvement. In fact, with the overhead of creating the commands, the efficiency may actually have deteriorated. The real power of GPU-driven rendering is in dynamic culling and level of detail. When you combine the technique of creating a command list on the GPU with other techniques such as mesh shading, you’ll realize the full power of the GPU. The next chapter will introduce you to the mesh shading pipeline.

Key Points

• You can create commands in indirect command buffers on either the CPU or the GPU.

• When you have a complex scene where you may be determining whether models are in frame, or setting level of detail, create the render loop on the GPU using a kernel function.

• You can use argument buffers to create large blocks of memory containing an entire scene. Use the GPU address of buffers, held in unified memory, as pointers in the argument buffer.

Where to Go From Here?

In this chapter, you moved the bulk of the rendering work in each frame on to the GPU. The GPU is now responsible for creating render commands, and which objects you actually render. Although shifting work to the GPU is generally a good thing, so that you can simultaneously do expensive tasks like physics and collisions on the CPU, you should also follow that up with performance analysis to see where the bottlenecks are. You can read more about this in Chapter 30, “Profiling”.

GPU-driven rendering is a fairly recent concept, and the best resources are Apple’s WWDC sessions listed in references.markdown in the resources folder for this chapter.

Apple’s sample: Modern Rendering with Metal renders Amazon’s huge Bistro scene, with heavy mesh resources and many lights. The sample uses several advanced techniques including indirect command buffers for GPU-driven rendering, and is the best project to tear apart and analyze.