Dispatch Geometry

A Metal dispatch launches a grid of threadgroups of threads: CUDA's grid/block/thread hierarchy under different names, with two API-level traps for the CUDA-trained.

The hierarchy, top to bottom:

The three levels, and both traps in one frame: the two dispatch calls count different units (boxes vs threads, with the ragged partials dispatchThreads quietly creates), and the threadgroup's size is promised to the compiler at pipeline time, not launch time.

Trap one: two dispatch calls with different units. dispatchThreadgroups(gridSize, threadsPerThreadgroup:) counts the grid in threadgroups, exactly CUDA's <<<numBlocks, blockDim>>>. But dispatchThreads(gridSize, threadsPerThreadgroup:) counts the grid in threads, and quietly handles ragged edges by making partial threadgroups (no guard if (i < n) needed, but also no guarantee of full simdgroups, which breaks kernels that assume cooperative full-width loads). Framework code uses both: MLX's custom-kernel API exposes grid in threads. When a ported kernel reads garbage at the edges or a cooperative load goes wrong, check which convention the dispatching layer uses before checking anything else.

Trap two: threadgroup size is a pipeline-time contract, not a launch-time detail. The compiler allocates registers when it builds the pipeline state, before it knows your launch dimensions, so it assumes the maximum unless the kernel promises otherwise with max_total_threads_per_threadgroup. In CUDA, __launch_bounds__ is a tuning hint; here its Metal twin is routinely worth integer factors and every production kernel in the case studies carries it.

Sizing instincts that transfer: threadgroup sizes are multiples of 32; 128 threads (4 simdgroups) is the workhorse size in every GEMM and attention kernel this glossary reads; occupancy saturates at lower residency than on an SM, so you rarely chase giant threadgroups. Grid-stride loops exist here too but are rarer, since dispatches are cheap to size exactly and non-uniform threadgroups cover the ragged case.

The index-attribute vocabulary ([[thread_position_in_grid]] and friends) replaces CUDA's blockIdx * blockDim + threadIdx arithmetic. Ask for the composed index directly rather than computing it.