Hello Metal Example

The 'each thread writes its coordinates to a buffer' pattern

Metal doesn't have GPU printf. Instead, the convention is: each thread writes its identity into a device buffer, and after the command buffer completes, the CPU reads the buffer and prints. Verbose, but it forces the round-trip you'll need for every real workload anyway.

Two files: hello.metal (the kernel) and hello_host.swift (the launch driver + result printer). Build with metal + metallib + swiftc, run as a single binary.

Code

hello.metal — the kernel·metal

#include <metal_stdlib>
using namespace metal;

kernel void hello(
    device uint2 *out               [[buffer(0)]],
    uint  tid_in_tg                 [[thread_index_in_threadgroup]],
    uint  tg_id_in_g                [[threadgroup_position_in_grid]])
{
    const uint threadsPerTG = 4;
    uint gid = tg_id_in_g * threadsPerTG + tid_in_tg;
    out[gid] = uint2(tg_id_in_g, tid_in_tg);
}

hello_host.swift — Swift driver·swift

import Metal
import Foundation

func main() throws {
    guard let device = MTLCreateSystemDefaultDevice() else {
        fatalError("No Metal-capable GPU")
    }
    let lib = try device.makeLibrary(URL: URL(fileURLWithPath: "hello.metallib"))
    let fn = lib.makeFunction(name: "hello")!
    let pipeline = try device.makeComputePipelineState(function: fn)

    let threadsPerTG = 4
    let threadgroups = 1
    let total = threadsPerTG * threadgroups

    let outBuf = device.makeBuffer(
        length: total * MemoryLayout<SIMD2<UInt32>>.stride,
        options: .storageModeShared)!

    let queue = device.makeCommandQueue()!
    let cb = queue.makeCommandBuffer()!
    let enc = cb.makeComputeCommandEncoder()!
    enc.setComputePipelineState(pipeline)
    enc.setBuffer(outBuf, offset: 0, index: 0)
    enc.dispatchThreadgroups(
        MTLSize(width: threadgroups, height: 1, depth: 1),
        threadsPerThreadgroup: MTLSize(width: threadsPerTG, height: 1, depth: 1))
    enc.endEncoding()
    cb.commit(); cb.waitUntilCompleted()

    let p = outBuf.contents().bindMemory(to: SIMD2<UInt32>.self, capacity: total)
    for i in 0..<total {
        print("Hello from threadgroup \(p[i].x), thread \(p[i].y)")
    }
}

do { try main() } catch { print("Error: \(error)"); exit(1) }

Build + run·bash

mkdir -p build
xcrun -sdk macosx metal -c hello.metal -o build/hello.air
xcrun -sdk macosx metallib build/hello.air -o build/hello.metallib
xcrun -sdk macosx swiftc hello_host.swift \
    -framework Metal -framework Foundation \
    -o build/hello_host

cd build && ./hello_host
# Hello from threadgroup 0, thread 0
# Hello from threadgroup 0, thread 1
# Hello from threadgroup 0, thread 2
# Hello from threadgroup 0, thread 3

Exercise

Save both files, run the three build commands, and confirm the four Hello lines print. Then change threadgroups = 1 to threadgroups = 2 in hello_host.swift AND set threadsPerTG to 8 in hello.metal (yes, both — they have to agree). Rebuild, run. You should see 16 lines, possibly out of order. That mismatch hint is intentional: in Metal, the kernel's hard-coded threadsPerTG and the host's threadsPerThreadgroup dispatch arg must match, and forgetting the kernel side is one of the top three Metal bugs.

The 'each thread writes its coordinates to a buffer' pattern

Code

External links

Exercise

Progress

Comments 0