metalworkingGitHub ↗
/machine/gpu-core

GPU Core

The GPU core is the M-series GPU's unit of compute: it owns a register file, threadgroup memory, and ALU pipes, and schedules simdgroups onto them.

CUDA equivalent: the Streaming Multiprocessor↗. The mapping is close enough that you can reuse most SM intuition directly.

Chip tiers differ only in how many cores you get and how much bandwidth feeds them: a base M5 has 10 cores at ~153 GB/s; an M5 Max has 40 cores at ~614 GB/s. There is no Apple datacenter part. The top of the line is a laptop/desktop chip, and everything about the design follows from that: Apple runs cores low-clocked and wide, tuning for performance per watt rather than peak throughput. An H100 will crush any M-series chip on raw FLOPs; the M-series counter is unified memory capacity and efficiency.

NVIDIA — A100 APPLE — M-SERIES SM 0 REGISTERS 256 KB per SM L1 / SHARED MEMORY 192 KB, configurable SM 107 108 SMs in A100 REGISTERS 256 KB per SM L1 / SHARED MEMORY 192 KB, configurable GPU CORE 0 REGISTERS ~208 KB per core THREADGROUP+TILE 32 KB / threadgroup L1 8 KB D · 12 KB I GPU CORE N−1 10 in M5 · 40 in M5 Max REGISTERS ~208 KB per core THREADGROUP+TILE 32 KB / threadgroup L1 8 KB D · 12 KB I L2 CACHE — 40 MB GPU LAST-LEVEL CACHE (L2) — 256 KB…1.5 MB BY CHIP ( no equivalent tier ) SYSTEM LEVEL CACHE (SLC) — 48 MB IN M1 MAX · SHARED WITH CPU + ANE GLOBAL MEMORY — HBM, 40 GB, ~1.6 TB/s GPU-only: the CPU cannot see this pool UNIFIED MEMORY — LPDDR, ONE POOL FOR CPU + GPU + ANE M5: ~153 GB/s · M5 Max: ~614 GB/s PCIe ~32 GB/s · cudaMemcpy CPU + HOST RAM — A SEPARATE POOL M3 AND LATER — DYNAMIC CACHING: registers, threadgroup, tile and stack come from one on-core pool

The two hierarchies side by side. Reading down: Apple inverts the register/shared-memory ratio, has no 40 MB L2 safety net, and adds an SLC tier. It then wins the bottom row, where NVIDIA's DRAM is GPU-only with the CPU across PCIe and Apple's is one shared pool. On M3 and later, Dynamic Caching allocates the per-core memories on demand from one cache pool instead of fixed partitions.

Inside a core, the numbers that matter:

  • Register file: ~208 KB, nearly SM-sized, and the core's real working memory.
  • Threadgroup memory: 32 KB per threadgroup, much smaller than CUDA's shared memory. The register/shared ratio is inverted relative to what you're used to, and that inversion drives kernel design here.
  • Caches are tiny: ~8 KB L1 data, ~12 KB instruction per core. Apple spent the transistors on registers instead. Plan on explicit reuse, not cache locality.
  • ALU saturation at ~24 resident simdgroups (~768 threads); see occupancy. You don't need SM-style 2048-thread residency to fill the machine.

The ALU pipes prefer 16-bit: F16 issues with measurably shorter dependent-instruction stalls than F32. Add the halved register pressure and you get the local wisdom, "use F16 everywhere". Matrix work runs through simdgroup_matrix, an 8×8 tile-multiply primitive closer to a fast wide-FMA arrangement than a separate tensor-core unit.

None of these numbers come from Apple. The microarchitecture was reverse engineered by the community, chiefly in philipturner/metal-benchmarks↗ (MIT), which is this section's primary source and the reference to return to whenever a kernel underperforms for no visible reason.