metalworkingGitHub ↗
/machine/nvidia-to-metal

NVIDIA to Metal

Every NVIDIA and CUDA term next to its Apple counterpart, on one page: the chip, the kernel code, the toolchain, and the libraries. Each row says whether your intuition transfers and links to the page that has the detail.

Read this page first, then use it as the index for the rest of the glossary. The other pages each explain one row.

The levels

Two hierarchies exist at the same time. One is software: the threads that one kernel launch creates. The other is hardware: the parts of the chip that run them. Most confusion comes from mixing the two, so keep them in separate columns.

SOFTWARE: WHAT ONE LAUNCH CREATES HARDWARE: WHAT IS ON THE CHIP THREAD one run of your kernel function, with its own registers NVIDIA: CUDA thread ARITHMETIC UNIT executes the instructions of a thread NVIDIA: CUDA core runs on THREADGROUP you choose its size: 128 threads here threadgroup memory: 32 KB, shared by all its threads simdgroup 0 simdgroup 1 simdgroup 2 simdgroup 3 simdgroup: 32 threads that execute the same instruction together NVIDIA: thread block with shared memory; a simdgroup is a warp GPU CORE holds several threadgroups at one time register file ~208 KB threadgroup memory ~60 KB L1 data cache ~8 KB bar length is proportional to size arithmetic units NVIDIA: Streaming Multiprocessor (SM) name trap: an NVIDIA "CUDA core" is one arithmetic unit, not this box runs on one core GRID all the threadgroups of one launch GPU 10 cores on M5, 40 on M5 Max threadgroup threadgroup threadgroup threadgroup threadgroup ... GPU core GPU core GPU core GPU core GPU core ... device memory (MTLBuffer): every thread can read and write it unified memory: the same memory the CPU uses, so no copy NVIDIA: kernel grid with global memory NVIDIA: GPU with GPU RAM on the card; the host must copy data to it is spread over all the cores

Left: what one launch creates. Right: the hardware that runs it. Each row reads across. The NVIDIA name for each box is under it.

How the two columns connect:

  • A threadgroup runs on one core. Its threads can share threadgroup memory because that memory is inside the core.
  • A core holds several threadgroups at one time. They divide its registers and threadgroup memory between them, and the number that fit is the occupancy.
  • The core schedules simdgroups, not single threads. About 24 simdgroups (about 768 threads) keep all its arithmetic units busy.
  • A grid is usually larger than the chip. The threadgroups that do not fit wait, and each one starts when a core has room.
  • "Core" names a different level on each side. An Apple GPU core is the size of an NVIDIA Streaming Multiprocessor. An NVIDIA "CUDA core" is one arithmetic unit inside it.

If you can change it in your code, it is software: the grid, the threadgroup size, the number of simdgroups. If you can change it only with a different chip, it is hardware.

The chip

The last column uses three words:

  • Transfers: the concept and the mechanism match. Reuse what you know.
  • Resized: the same concept with different numbers, and the numbers change which kernel design wins.
  • Unlearn: the CUDA habit is wrong here, because the feature is absent or its cost has the opposite sign.
What it is NVIDIA Apple M-series Verdict
Unit of compute Streaming Multiprocessor (SM) GPU core Transfers
32 threads that run in lockstep warp simdgroup Transfers
Register file ~256 KB per SM ~208 KB per core Resized: similar size, but here it is the main working memory
Scratch memory shared by a block shared memory, up to 228 KB per SM, configurable threadgroup memory, 32 KB per threadgroup, fixed Resized: a staging buffer, not the place where tiles stay
Caches 40 MB L2 (A100) ~8 KB L1 data per core, a small L2, and an SLC tier Resized: plan on explicit reuse, not cache locality
Main memory HBM on the card at multiple TB/s; host RAM is across PCIe unified memory shared with the CPU, ~153 GB/s (M5) to ~614 GB/s (M5 Max) Unlearn: no transfers to manage, but bandwidth is the scarce resource
Threads needed to fill the ALUs 2048 resident threads per SM ~24 simdgroups (~768 threads) per core Resized
Matrix unit Tensor Core none before M5, where simdgroup_matrix runs on the regular FP32 pipes; neural accelerators on M5 Unlearn before M5
Reason to use 16-bit floats unlocks Tensor Core throughput shorter stalls and half the register pressure Resized: same advice, different reason
Fast exponent the SFU's MUFU.EX2 fast::exp2 Transfers
Float atomics atomicAdd(float*) is a cheap hardware instruction emulated Unlearn
Copy engine that overlaps with compute cp.async / TMA in the hardware but not exposed; the Metal 4 compiler rejects it Unlearn
Cost of a block-wide barrier high enough that avoiding __syncthreads() is an optimization ~2 cycles Unlearn: barriers are nearly free
Largest part datacenter GPU (H100) laptop and desktop chip (M5 Max, 40 cores); there is no datacenter part An H100 wins on raw FLOPs; Apple wins on memory capacity and efficiency
More than one GPU NVLink inside one chassis one device per machine; Thunderbolt 5 RDMA between machines Unlearn

None of the Apple numbers come from Apple. The community measured them, chiefly in philipturner/metal-benchmarks↗.

The kernel code

CUDA C++ and the Metal Shading Language (MSL) are close enough that you can read one if you can read the other.

CUDA Metal Detail
grid, thread block, warp, thread grid, threadgroup, simdgroup, thread Dispatch geometry
__global__ void f(...) kernel void f(...) MSL
threadIdx, blockIdx, blockDim attribute-tagged parameters such as [[thread_position_in_threadgroup]] MSL
__shared__ float s[N]; threadgroup float s[N]; Threadgroup memory
__launch_bounds__(n) __attribute__((max_total_threads_per_threadgroup(n))) Registers
__syncthreads() threadgroup_barrier(mem_flags::mem_threadgroup) Synchronization
__syncwarp() simdgroup_barrier(mem_flags::mem_none) Synchronization
wmma fragments, mma.sync simdgroup_matrix, in 8×8 tiles simdgroup_matrix
cp.async no supported equivalent simdgroup_async_copy
TMA plus wmma descriptors MTLTensor and MPP (Metal 4) MTLTensor and MPP
one compiled kernel per configuration (template instantiation, NVRTC) one kernel with function constants, bound when the pipeline is built Function constants

The toolchain and runtime

CUDA Metal Detail
device and context MTLDevice, usually one per chip Metal, the API
stream MTLCommandQueue Command buffers
kernel launch <<<...>>> a dispatch encoded into a MTLCommandBuffer, then committed Command buffers
cudaMalloc, then cudaMemcpy from the host MTLBuffer; the CPU and GPU see the same bytes Unified memory
nvcc the metal frontend (xcrun metal), or compile at runtime Compilation pipeline
PTX AIR, Apple's portable intermediate representation Compilation pipeline
fatbin metallib Compilation pipeline
cubin, SASS the GPU binary inside a pipeline state, always built on the user's machine Compilation pipeline
Nsight Compute nothing comparable Profiling
cuobjdump, nvdisasm applegpu, a community disassembler Disassembly

The libraries

NVIDIA stack Apple stack Detail
PyTorch, cuBLAS/cuDNN, and parts of Triton MLX, one codebase MLX, an overview
cuBLAS, cuDNN MPS, which open kernels have matched or beaten MPS
CUTLASS steel Steel
cuDNN fused attention, the FlashAttention library mx.fast mx.fast
Triton, for custom kernels written from Python mx.fast.metal_kernel mx.fast
torch.compile mx.compile, which does elementwise fusion and stops there mx.compile
PyTorch's asynchronous dispatch lazy evaluation: nothing runs until mx.eval Lazy evaluation
weight-only quantization kernels (TensorRT-LLM, AWQ, GPTQ) MLX quantization Quantization
NCCL over NVLink mx.distributed over Thunderbolt 5 RDMA Distributed

What to unlearn

The rows marked Unlearn and Resized become these habits:

  1. Keep the working set in registers, not in shared memory. Threadgroup memory is 32 KB, so tiles pass through it and the results stay in registers. The failure is a spill, which ran 10× slower in the measured case.
  2. Use barriers freely. A barrier costs ~2 cycles, so choose the algorithm with the cleaner staging pattern even if it syncs more often.
  3. Do not accumulate with float atomics. They are emulated, and kernels that depend on them need a different structure.
  4. Do not plan on copy and compute overlap. Without cp.async, tiles move by cooperative loads, and double buffering gets its overlap from instruction-level parallelism.
  5. Expect to be bandwidth bound. More kernels are limited by memory bandwidth than CUDA experience suggests, so arithmetic intensity decides most designs, and fusion and quantization pay more here.
  6. Batch dispatches and sync rarely. A launch is not a function call. Put many dispatches in one command buffer and wait for the GPU as rarely as possible.
  7. Build your own measurements. There is no Nsight Compute, so you assemble the roofline yourself and change one thing at a time.

What transfers without change: warp intuition, the occupancy model, tiling, and the algorithms themselves, such as online softmax and flash attention.

Edit this page on GitHub ↗