Metal, MPS and Tensor APIs: The GPU's Public Doors

"There is exactly one way onto Apple's GPU, and everything else is a door built in front of it."

The One Way In

Every program that runs on an Apple GPU does so through Metal: Apple's graphics and compute API, the only one the hardware exposes. A Metal compute kernel is a small function in the Metal Shading Language, compiled by Apple's compiler, dispatched over a grid of threads by a command queue. That is the door. There is no CUDA, no OpenCL that Apple still maintains, no Vulkan from Apple. When a framework says it "supports Apple silicon", it means it has written or borrowed Metal kernels for the operations it needs. The code block writes one of those kernels through MLX's fast.metal_kernel, which lets you hand Metal Shading Language source to the GPU from Python and get an array back.

The Doors Built in Front of It

Door	What it is	Who walks through it	Evidence
Metal Performance Shaders (MPS)	Apple's library of pre-written Metal kernels — matrix multiply, convolution, image ops	Core ML underneath; PyTorch's backend, which took its name from it	vendor (Apple docs)
MLX	Apple's array framework: its own Metal kernels (and a CPU backend), lazy evaluation, unified-memory-native	mlx-lm, Ollama's MLX engine, this quest's lab, most of the household's local inference	vendor (MLX docs)
PyTorch MPS backend	PyTorch tensors placed on `'mps'`, dispatched to Metal (originally via MPS, now largely custom kernels)	Hugging Face pipelines, diffusers, anything written for CUDA first — the household's image-generation engine	vendor (PyTorch docs)
llama.cpp's Metal backend	Hand-written Metal kernels for quantized inference	llama.cpp, LM Studio, Ollama's original engine	vendor (llama.cpp README: "Apple silicon is a first-class citizen — optimized via ARM NEON, Accelerate and Metal frameworks")
Core ML	A compiled-model runtime that chooses CPU, GPU or Neural Engine per layer	App-embedded models; the household's on-device speech recognizer	vendor (Apple docs)
Metal 4 Tensor APIs	`MTLTensor` and tensor operations in Metal 4 (macOS 26+); the way to program the M5 GPU's Neural Accelerators directly	MLX from 0.30 ("Support for Neural Accelerators on M5 (macOS >= 26.2)")	vendor (Apple, MLX release notes)

The table has a shape worth noticing. The doors that dominate local language-model inference — MLX, llama.cpp — are the ones with hand-written kernels for quantized matrix-vector products, because decode is bandwidth-bound and a kernel that reads each weight exactly once at full width is the whole game. The doors that dominate everything else — PyTorch, Core ML — are the ones that inherited a programming model from somewhere else and pay a translation cost for it. Ollama's 2026 move is the clearest sign of where the wind blows: "Ollama is now powered by MLX on Apple Silicon in preview", with llama.cpp kept alongside.

The M5 Door, Stated as a Claim

With the M5 generation Apple put a Neural Accelerator in every GPU core and opened it through Metal 4: "Developers can also build solutions for their apps by directly programming the Neural Accelerators using Tensor APIs in Metal 4." MLX's own post says it "leverages the Tensor Operations (TensorOps) and Metal Performance Primitives framework introduced with Metal 4", and requires macOS 26.2 or later for it. Apple's numbers for the effect are the prefill numbers from track one — "up to 4x speedup compared to a M4 baseline for time-to-first-token" — and, for decode, a "19-27% performance boost … thanks to its greater memory bandwidth". Read that second clause carefully: Apple itself attributes the decode gain to bandwidth, not to the accelerators. Every M5 statement in this lesson is a vendor claim; this quest measures no GPU past M3 (the M5 Max laptop in the house sits under the plan's no-M5-measurement ruling), and the accelerators are a door this quest has not walked through.

Code

metal_door.py — a Metal Shading Language kernel dispatched from Python through MLX·python

#!/usr/bin/env python3
"""The one real door: a Metal compute kernel. MLX compiles the MSL body below
into a kernel, dispatches it over a grid, and hands back an MLX array in
unified memory. Requires mlx on Apple silicon."""
import time
import mlx.core as mx

SOURCE = """
    uint i = thread_position_in_grid.x;      // one GPU thread per element
    out[i] = 2.0f * inp[i] + 1.0f;           // the whole kernel: an affine map
"""
affine = mx.fast.metal_kernel(name="affine", input_names=["inp"], output_names=["out"], source=SOURCE)


def run(x: mx.array) -> mx.array:
    return affine(inputs=[x], grid=(x.size, 1, 1), threadgroup=(256, 1, 1),
                  output_shapes=[x.shape], output_dtypes=[mx.float32])[0]


x = mx.arange(8, dtype=mx.float32)
print(run(x))                                   # array([1, 3, 5, ..., 11, 13, 15])

n = 1 << 28                                      # 1 GiB in, 1 GiB out
big = mx.ones((n,), dtype=mx.float32); mx.eval(big)
mx.eval(run(big))                                # compile + warm
best = 1e9
for _ in range(3):
    t = time.perf_counter(); mx.eval(run(big)); best = min(best, time.perf_counter() - t)
print(f"custom Metal kernel, 1 GiB in + 1 GiB out: {best*1e3:.1f} ms -> {2*n*4/best/1e9:.0f} GB/s")
# office, M3 Ultra, mlx 0.32.2, 2026-09-15: 10.0 ms -> 215 GB/s (a naive one-element-per-thread kernel;
# MLX's own add reached 635 GB/s in stream.py — writing a fast kernel is its own craft)

Exercise

Run metal_door.py, then change the kernel body to compute the elementwise product of two inputs (add a second input name) and confirm the result on a small array. Then list every framework installed on your Mac that reaches the GPU and, for each, name the door from the table it uses. Which of them would let you write the kernel you just wrote?

Hint

Only MLX (fast.metal_kernel) and raw Metal let you write your own kernel from Python or Swift. PyTorch MPS and Core ML choose kernels for you; llama.cpp lets you edit its Metal source but not inject at runtime. The door you can walk through with your own code is the one worth knowing.