Metal Compute on Apple Silicon: Argument Buffers, Shared Memory, and the MPS Decision

The hidden cost is not the shader

You write a Metal compute kernel, profile it, and find the GPU is idle half the time. The shader math is trivial; the real bottleneck is the CPU encoding dozens of individual buffer and texture bindings into every command encoder. On Apple Silicon, Unified Memory Architecture (UMA) means the CPU and GPU share one physical memory pool, so the old "copy data to GPU" tax is gone. But the API call overhead and the discipline required to order access across that shared pool remain very real. The engineering decision that matters most is usually not whether to use Metal, but how you structure resource bindings and synchronization.

Unified Memory Architecture removes copies, not hazards

Because the CPU and GPU address the same physical memory, a MTLBuffer created with MTLResourceStorageModeShared can be read and written from both sides without an explicit blit or copy. That sounds like synchronization becomes optional. It does not. Shared storage mode only removes the transfer step; it does not remove the need to order operations. If a compute kernel writes to a buffer and a subsequent render pass reads it, you still need a barrier—either an MTLFence within a single command buffer or an MTLEvent across multiple buffers. Skipping that ordering produces race conditions that may manifest as torn reads or, worse, GPU hangs that are hard to reproduce.

Argument Buffers trade flexibility for fewer API calls

By default, binding resources means calling setBuffer:offset:atIndex: and setTexture:atIndex: for every texture and buffer a kernel needs. That is fine for a handful of resources, but it does not scale. An Argument Buffer packs many resources into a single MTLBuffer that the shader indexes as a structured argument. On the CPU side, you encode the arguments once with an MTLArgumentEncoder and bind the single buffer. The trade-off is indirection: the shader accesses resources through the argument buffer structure rather than direct binding slots, and you must explicitly tell the encoder about resource residency with useResource:usage: or useHeap:. If you skip residency, Metal cannot track lifetime and you may access freed memory.

A minimal compute dispatch pattern

The following pattern illustrates a typical compute pass. It is not a benchmark; treat it as a structural reference and verify against your target SDK. The .metal source runs on the GPU after compilation, while the Swift snippet runs on the CPU inside an Xcode project with the Metal framework linked.

In a .metal file:

#include <metal_stdlib>
using namespace metal;

kernel void scale_values(
    device float *values [[ buffer(0) ]],
    constant float &factor [[ buffer(1) ]],
    uint index [[ thread_position_in_grid ]]
) {
    values[index] *= factor;
}

On the host side, inside a view controller or command queue delegate:

guard let commandBuffer = commandQueue.makeCommandBuffer(),
      let encoder = commandBuffer.makeComputeCommandEncoder() else { return }

encoder.setComputePipelineState(computePipelineState)
encoder.setBuffer(sharedBuffer, offset: 0, index: 0)
var factor: Float = 2.0
encoder.setBytes(&factor, length: MemoryLayout<Float>.size, index: 1)

let gridSize = MTLSize(width: elementCount, height: 1, depth: 1)
let threadgroupSize = MTLSize(width: 256, height: 1, depth: 1)
encoder.dispatchThreads(gridSize, threadsPerThreadgroup: threadgroupSize)
encoder.endEncoding()
commandBuffer.commit()

Here, sharedBuffer uses shared storage mode so the CPU can initialize the data without a separate staging buffer. The threadgroupSize of 256 is a common starting point for simple kernels, but the optimal value depends on the device and the kernel's register usage. After commit(), the CPU can continue encoding other work; it does not wait for the GPU unless you call waitUntilCompleted().

MPS, custom MSL, and the profiling decision

Metal Performance Shaders (MPS) provides pre-optimized implementations for common operations—Gaussian blurs, matrix multiplication, and neural network layers among them. The instinct to write a custom MSL kernel for a standard image filter is often a false economy. MPS kernels are tuned for Apple's GPU architectures and receive updates with the OS. A practical decision process is:

1. If the operation is a standard primitive (convolution, blur, GEMM), prototype with MPS first.
2. If MPS lacks the exact operation or you need fused steps that MPS cannot express, write custom MSL.
3. Profile both paths with the Metal Debugger before committing.

The comparison is not always about raw shader speed. MPS may win because it reduces your CPU encoding burden and avoids custom synchronization bugs.

Verification and limitations

Two limitations deserve attention. First, MSL is proprietary to Apple platforms. Code is not portable to Vulkan or DirectX without a translation layer. If cross-platform support is a requirement, that cost must be factored in. Second, memory management on Apple Silicon is not unlimited. Allocating large MTLHeap objects or retaining many shared buffers can trigger system-level memory pressure. The OS may terminate your process with little warning. Monitor memory in Instruments and prefer heaps for transient, same-frame allocations rather than long-lived resource pools.

To verify your pipeline, use Xcode's Metal Debugger. Capture a frame and inspect the GPU timeline. Look for gaps between command buffers—those indicate CPU encoding bottlenecks that Argument Buffers can reduce. Check the dependency viewer to confirm your fences and events are actually serializing access. If a compute pass reads data the CPU just wrote, the debugger will show whether the hazard is tracked correctly.

Closing checklist

• Start with shared storage mode for data that both CPU and GPU touch, but never skip MTLFence or MTLEvent when ordering access.
• Group frequently bound resources into Argument Buffers once encoder CPU time shows up in the Metal Debugger.
• Reach for MPS before writing custom MSL for standard operations; keep custom kernels for logic MPS cannot express.
• Profile memory pressure in Instruments, especially if using heaps or large persistent buffers.

Because Metal APIs evolve, confirm the exact method signatures and storage mode availability against the Apple developer documentation for your current Xcode and OS version before shipping.