What are the steps of a Metal compute pass, what runs when, and what can you verify without a GPU frame?

18.8.3 · L5 · Interview

Predict first

Does the compute pass give the right answer, and when do errors appear?

import Foundation
import Metal

let source = """
#include <metal_stdlib>
using namespace metal;
kernel void addArrays(device const float* a [[buffer(0)]], device const float* b [[buffer(1)]],
                      device float* out [[buffer(2)]], uint i [[thread_position_in_grid]]) {
    out[i] = a[i] + b[i];
}
"""
guard let device = MTLCreateSystemDefaultDevice() else { fatalError("no Metal device") }
let n = 1_000_003                                  // deliberately not a multiple of the threadgroup size
let pipeline = try device.makeComputePipelineState(function: try device.makeLibrary(source: source, options: nil).makeFunction(name: "addArrays")!)
let a = (0..<n).map { Float($0) }, b = (0..<n).map { Float($0) * 2 }
let bufA = device.makeBuffer(bytes: a, length: n * 4, options: .storageModeShared)!
let bufB = device.makeBuffer(bytes: b, length: n * 4, options: .storageModeShared)!
let bufOut = device.makeBuffer(length: n * 4, options: .storageModeShared)!

let commands = device.makeCommandQueue()!.makeCommandBuffer()!
let encoder = commands.makeComputeCommandEncoder()!
encoder.setComputePipelineState(pipeline)
for (i, buf) in [bufA, bufB, bufOut].enumerated() { encoder.setBuffer(buf, offset: 0, index: i) }
let group = min(pipeline.maxTotalThreadsPerThreadgroup, 256)
encoder.dispatchThreads(MTLSize(width: n, height: 1, depth: 1), threadsPerThreadgroup: MTLSize(width: group, height: 1, depth: 1))
encoder.endEncoding()
print("before commit, status is notEnqueued:", commands.status == .notEnqueued)
commands.commit()
commands.waitUntilCompleted()
print("completed:", commands.status == .completed, "| error:", commands.error as Any)
let out = bufOut.contents().bindMemory(to: Float.self, capacity: n)
print("out[10] =", out[10], "| out[n-1] =", out[n - 1], "| expected last =", Float(n - 1) * 3)
print("a thread-execution width divides the group size:", group % pipeline.threadExecutionWidth == 0)
print("a missing kernel name gives nil:", try device.makeLibrary(source: source, options: nil).makeFunction(name: "nope") == nil)
do { _ = try device.makeLibrary(source: "kernel void broken( {", options: nil) } catch { print("a Metal source error throws at library build time") }

notEnqueued before commit, completed after wait, correct last element, nil for a missing kernel, throw for broken source
results are readable before commit
a missing kernel name crashes

Mental model. Writing instructions on a work order, handing it in with commit, and collecting the finished parts later; the order is only a plan until it is handed in.

The chain is: MTLDevice (nil means no Metal), a library (from .metal files in the bundle, or makeLibrary(source:) at run time as in the check), a function, a MTLComputePipelineState, buffers, a MTLCommandQueue, a command buffer, a compute encoder, dispatchThreads, endEncoding, commit. Nothing runs until commit: the check shows the command buffer is notEnqueued before it, then waits (waitUntilCompleted) and reads results through contents() of shared buffers. It added two million-element arrays (1,000,003 elements, deliberately not a multiple of the threadgroup size) and got the right last element, because dispatchThreads handles non-uniform grids. Errors split by time: a source error throws when you build the library, a wrong kernel name returns nil, and GPU execution errors appear on commandBuffer.error after completion.

In production do not block: add a completion handler or await the buffer, reuse the queue and pipeline states (building a pipeline is the expensive step), size threadgroups from the pipeline’s threadExecutionWidth and maxTotalThreadsPerThreadgroup instead of hard-coding, and mind storage modes (shared is CPU-visible; private is GPU-only and faster for data the CPU never reads). Check MTLCreateSystemDefaultDevice() for nil and supportsFamily before using features. Not run: iPhone hardware (this ran on a Mac GPU), render passes, and profiling, so no timings are claimed.

The GPU works only after commit; errors surface at library build, function lookup and completion.

Common traps

Reading results before completion. The GPU may not have run yet.
Building pipelines per frame. Cache them.

They might follow up with

Does the simulator have Metal? It has a device, but it does not necessarily report the same feature families as hardware (Q398 covers choosing frameworks).

What are the steps of a Metal compute pass, what runs when, and what can you verify without a GPU frame?

Predict first

Common traps

They might follow up with

Related questions