Apple M5 (base · 10 GPU cores · 24 GB) · the hardware, measured · Part 4 of the measurement project · census of 2026-09-01
Where a tensor operation goes on this chip, how it is split across cores, partitions and simdgroups, and why prefill and decode land on opposite sides of the same ridge point.
Part I pins down the M5's hardware by direct measurement: the chip layout, the full memory hierarchy from lane registers to DRAM, the inside of a GPU core, and how a matrix multiply is physically split across it. Part II applies those numbers to LLM inference: why prefill and decode are different problems, which kernels run where, and which optimizations the measurements say will pay. Every figure is measured on one base M5 (ten GPU cores, 24 GB) or labeled when it is not.
| Part I · the hardware | Memory latency rises in four steps, at about 30, 54, 80 and 200 ns. Each GPU core has about 128 KiB of on-chip storage, and registers, threadgroup memory and L1 all draw from it. The chip-wide system level cache holds 4–8 MiB for GPU access. Each core carries a Neural Accelerator; together they deliver 15.4 TFLOPS fp16, and an ordinary Metal kernel can drive them bit-exact. |
| Part II · the LLM | 15.4 TFLOPS over 120.4 GB/s puts the ridge point at 128 FLOP per byte. Prefill clears it, so the Neural Accelerators limit it. Decode sits 32–128× below it, so memory bandwidth limits it. The ranked list of work at the end follows from these numbers. |
What the silicon is, established by probes on this machine.
The blocks a tensor passes through, and which of them a Metal kernel can reach. Then the two nonstandard blocks that make Apple silicon different, and which of them matters for LLMs.
CPU cores, GPU cores and a cache hierarchy are standard architecture; every laptop has them. Apple silicon adds two kinds of specialized machine-learning hardware that are not standard. One of them is the reason an M5 runs LLMs as fast as it does. The other cannot help this workload at all.
| Block | What it is | Why it matters here |
|---|---|---|
| NAX the matrix-multiply unit inside each GPU core |
A dedicated matrix unit inside each GPU core, new in the M5 generation, programmed from ordinary Metal shaders via Metal 4 tensor operations. Apple's public headers call it the Neural Accelerator (NA); Apple's MLX code calls it NAX. Both names mean the same silicon, and this page uses them interchangeably. | It sets the compute ceiling of the chip. 15.4 TFLOPS fp16, 2.5× what the shader ALUs can do, and a custom kernel reaches it bit-exact, verified here. Every prefill number in Part II rests on it. |
| ANE the Apple Neural Engine, a separate machine-learning block beside the GPU |
A fixed-function fp16 block beside the GPU, present in Apple chips since 2017. Core ML drives it in practice; third-party drivers reach it directly through a private, unsupported API. | Not usable for this workload. On the same part it has a tenth of the GPU's memory bandwidth and a third of its matmul throughput, so it loses at prefill and at decode; section 06 has the numbers. No MLX path and no bf16 keep it out of reach regardless. It appears again in sections 06, 10 and 11. |
The Neural Accelerators exist only from GPU architecture generation 17, which is the M5 and A19 family. Every earlier chip lacks them.
A Metal 4 kernel issues matrix work through the tensor operations (mpp::tensor_ops::matmul2d over metal::tensor operands). This project verified that path from mx.fast.metal_kernel: the result is bit-exact against mx.matmul, and MLX needs no rebuild. The accepted operand types are fp16, bf16, int8 and packed int4. The toolchain (build 17F109) has no fp8 type.
Every number on this page is a statement about one of six storage levels. A census (a fixed set of probes, run three times under one protocol) on 2026-09-01 gave every level a number measured on this machine. The few unknowns that remain are narrow, and they are listed at the end of this section.
census/hierarchy_probe.py, three rounds, 2026-09-01, results in telemetry/hierarchy.csv.| Level | Where it appears | What it decides |
|---|---|---|
| lane registers ×10 files | section 04, the cooperative tensor | the accumulator never leaves them across the K loop, which is why accumulation costs no memory traffic |
| threadgroup memory ×10 pools | section 04 tile calculator; section 07 hd=256 tile | 32 KB caps the staged tile, which caps arithmetic intensity, which caps attainable TFLOPS |
| L1 ×10 / L2 ×1 shared | the ~30 ns and ~54 ns rungs of the measured ladder | L1 holds 128–256 KiB per core; the next level holds 0.5–1 MiB as seen from one core. Streams from several threadgroups contend only mildly there |
| system level cache ×1 | section 08, the MoE gather lane | a gather inside 6 MB runs at 29.5 GB/s instead of 7.0; a MoE kernel that keeps its expert reads inside the cache recovers that 4× |
| DRAM ×1 | sections 05 and 08 | 120.4 GB/s is the decode ceiling and half of the ridge point |
Each probe ran three times on 2026-09-01, with a warm-up pass discarded and with no other GPU work on the machine. The probes are in census/hierarchy_probe.py; the results are in telemetry/hierarchy.csv.
| Question | Probe | Answer |
|---|---|---|
| latency at every level | single-thread pointer chase, 2 KiB → 128 MiB | ~30 / ~54 / ~80 / ~200 ns ladder with knees at 128–256 KiB, 0.5–1 MiB, 4–8 MiB |
| threadgroup memory against device memory | identical chase staged in threadgroup memory | 29.85 ns, equal to L1; strided reads 617 vs 567 GB/s aggregate, a 9% edge |
| physical pool per core | fixed work vs allocation, 1→32 KiB | ≈128 KiB; flat to 8 KiB, then steps matching floor(128/alloc) |
| register pressure | independent FMA streams, 4→160 live floats/thread | clean to 96 floats/thread; 1.35× cost at 128, 2.7× at 160: a knee, not a file size |
| shared-level contention | private streams × 1→20 threadgroups | flat at one 128 KiB footprint per core; halves at two; 512 KiB footprints degrade 11.4 → 7.8 GB/s per threadgroup by 10 |
| SLC capacity against TLB reach | full-random vs page-local chase at equal footprint | the 4→8 MiB cliff is capacity: both curves climb it. Translation adds +11 ns at 64 MiB, +47 ns at 128 MiB |
Each chart plots the three-round medians from telemetry/hierarchy.csv. Above each chart: what it shows, and how the measurement was made. Below it: what to read off.
Still open after the census: whether the 0.5–1 MiB level is per-core or shared (a one-core chase cannot tell), the register file's absolute size (the knee is a bound; the compiler may spill before the file fills), and peak threadgroup memory bandwidth (the 617 GB/s is loop-limited, a floor).
MTLDevice.maxThreadgroupMemoryLength = 32,768 B · architecture = applegpu_g17g · supportsFamily(Apple10) = true · 24 GiB unified, 18,186 MiB recommended working setOne GPU core, not to be confused with a CPU core: they are separate blocks that share only the SLC and DRAM. Apple documents none of this. The structure below is the standard Apple GPU core arrangement, plus a datapath shape inferred by a third party from throughput on another part of this GPU generation. Only the aggregate figures are measured on M5.
Measured on M5, int8 reaches 18.6 TOPS against 15.4 TFLOPS fp16. Both come from this project's census. Packed int4 is an accepted operand type in the tensor-op API, and its rate has not been measured yet. Section 11 lists that measurement as lane 3.
C = A · B, with C of shape M × N. The decomposition is the same at every level: a horizontal band of A meets a vertical band of B, and their intersection is a tile of C. That shape repeats at three levels, and the only things that change are the tile sizes and where the operands live.
BM × BK piece of A and one BK × BN piece of B at a time and accumulating into the tile.Each staged piece of A is reused BN times, once per output column in the tile. Each staged piece of B is reused BM times. That reuse is what lifts arithmetic intensity above the ridge, and one line of arithmetic gives it.
FLOPs = 2 · BM · BN · K
bytes = 2 · K · (BM + BN)
AI = BM · BN / (BM + BN) → for a square tile, simply BM / 2
K cancels. Depth does not raise intensity; only the face of the tile does. To reach the measured ridge of 128 FLOP per byte a square tile must be 256 × 256. BK is then chosen to fit the 32 KB budget.
The staged tiles are subdivided among the simdgroups, and the simdgroup's share is subdivided again into the accelerator's native operation. Each level repeats the band-meets-band geometry, one storage class down.
The same numbers, applied: why prefill and decode behave differently, and what the measurements say to build.
This project’s census measured two ceilings independently. Their ratio is the ridge point, and it sorts every operation in an LLM into compute-bound or bandwidth-bound before any kernel is written.
An operation needs 128 FLOPs of arithmetic for every byte it moves to saturate the Neural Accelerators. Below that, it is waiting on memory and the NA idles no matter how the kernel is written.
Prefill clears the ridge. Decode is not close.
A batch-1 decode GEMV does two FLOPs per weight. At fp16 that is 2 FLOP / 2 B = 1.0 FLOP/byte. At 4-bit weights it is 2 FLOP / 0.5 B = 4.0 FLOP/byte. Against a ridge of 128, decode sits 32× to 128× below it.
Attainable compute at that intensity is 120.4 GB/s × 4 = 0.48 TFLOPS, which is 3.1% of the NA peak. During decode the Neural Accelerators are structurally idle. No kernel recovers that; the arithmetic intensity of the problem forbids it.
Prefill and decode, plotted against the ridge this project’s census measured.
Bryngelson's study of the Apple Neural Engine (arXiv:2606.22283) measures all three processors on an M5 part. The absolute roofs are about twice this project's base-M5 figures, consistent with a larger M5 configuration, but the GPU ridge agrees to within 5%, and the ridge is the number every conclusion in Part II rests on.
| Device | Compute roof | Bandwidth roof | Ridge |
|---|---|---|---|
| ANE measured, their part | 10.2 TFLOP/s matmul 18.8 conv | 24.1 GB/s standalone | 424 |
| GPU measured, their part | 30.9 TFLOP/s | 229.7 GB/s | 134 |
| CPU measured, their part | 1.9 TFLOP/s | 130.4 GB/s | 15 |
| GPU · this project, base M5 | 15.4 TFLOP/s | 120.4 GB/s | 128 |
Their GPU roofs are 2.00× and 1.91× this project's. Both scaled together, so the ridge barely moved. That is the signature of a bigger part rather than a different measurement convention, and it means the compute-versus-bandwidth boundary is a property of the M5 design rather than of one machine.
Where each measured operation lands on that roofline. Two metrics appear: the July figures count the full score rectangle as work, which credits a causal kernel for blocks it skips; from 2026-08-27 the page counts only visible work, which lowers every causal prefill percentage (hd=128 reads about 79% on it, not 87–89%).
| Operation | Arithmetic intensity | Bound by | Attainable | NA utilization |
|---|---|---|---|---|
| Prefill attention hd=128, causal, 4k–32k | > 128 | compute · the NA | ≈12.2 TFLOPS | ≈79% |
| Prefill attention hd=256, causal, upstream fused kernel (MLX #3842) | > 128 | compute · the NA, accumulator-bound | 8.9–9.5 TFLOPS | 57.5–61.5% |
| Prefill attention hd=256, MLX 0.32.0 unfused fallback (July 2026) | > 128 | compute · but untiled | 5.3–6.7 TFLOPS | 34–44% July metric |
| Decode attention GQA, 32k KV | < 8 | memory · KV streaming | — | 64% of the bandwidth roof |
| Decode GEMV int4 weights, B=1 | 4.0 | memory · weight streaming | 0.48 TFLOPS | 3.1% |
| Decode GEMV fp16 weights, B=1 | 1.0 | memory · weight streaming | 0.12 TFLOPS | 0.8% |
| MoE expert gather | < 1 | memory · random access | 7.0 GB/s floor | ≈0% |
Prefill parallelizes over output tiles and decode over the reduction, because different silicon limits each.
Prefill has thousands of query positions to work on at once. Tiles are large, staged bytes get reused many times over, arithmetic intensity clears 128, and the job is to keep the ten Neural Accelerators fed. Parallelism is over output tiles.
Decode has one query position. There is nothing to batch over, every weight is read once and used twice, and the job is to move bytes. Parallelism is over the reduction, split-K across the KV length, not over the output. The NA cannot help.
The Apple Neural Engine is slower than the GPU at prefill and at decode, and the reason is bandwidth. Bryngelson measures all three processors on one M5 part, so the comparison below is internally consistent even though that part is larger than this machine; the last paragraph of this subsection estimates what the same ratios are on the base M5.
Three ceilings decide it, and each ratio goes against the ANE:
| Ceiling | ANE | GPU | What the ratio means |
|---|---|---|---|
| Compute roof, matmul | 10.2 TFLOP/s | 30.9 TFLOP/s | The ANE does not win on arithmetic either. It has a third of the GPU's matmul throughput on the same chip. |
| Bandwidth roof | 24.1 GB/s | 229.7 GB/s | The decisive ratio. A factor of 9.5. Anything limited by memory traffic runs at about a tenth of the GPU's rate. |
| Ridge point | 424 FLOP/byte | 134 FLOP/byte | An operation needs 3.2× more arithmetic per byte to saturate the ANE than to saturate the GPU. |
Applying those ceilings at the intensities this workload runs at: decode's intensity is 1.0 FLOP per byte at fp16 weights and 4.0 at 4-bit, from section 05; prefill attention clears the GPU's ridge and is placed here at the 400 FLOP per byte the roofline chart above uses.
| Operation | Arithmetic intensity | ANE attains | GPU attains | GPU advantage |
|---|---|---|---|---|
| Decode GEMV fp16 weights | 1.0 | 0.02 TFLOP/s | 0.23 TFLOP/s | 9.5× |
| Decode GEMV int4 weights | 4.0 | 0.10 TFLOP/s | 0.92 TFLOP/s | 9.5× |
| Prefill attention long sequence | ≈400 | 9.6 TFLOP/s | 30.9 TFLOP/s | 3.2× |
The ANE does not reach its own compute roof even at prefill intensity. At 400 FLOP per byte it is still on its memory roof, attaining 9.6 of its 10.2 TFLOP/s, because its ridge sits at 424. The GPU at that same intensity is compute-bound and delivers its full 30.9. The gap is 9.5× at every intensity below the GPU's ridge, narrows across the span between the two ridges, and settles at 3.0×, the compute ratio, past 424. There is no intensity at which the ANE wins.
The mechanism behind the bandwidth number is a working-set limit. Bryngelson finds the ANE's speed advantage stops once the operand working set passes 2 MB. Inside that window it is fast. Past it, everything streams at 24.1 GB/s. Decode's working set is the weights, gigabytes even at 4-bit, plus a KV cache that reaches hundreds of megabytes at 32k context: two to four orders of magnitude past the window. That is also why long-sequence attention and large square matmul, which are this project's prefill shapes, both measure faster on the GPU on their part.
What survives is energy, not throughput. The same study measures the ANE about 13× more energy efficient than the GPU on M5. Nothing above contradicts that: a block with a tenth of the bandwidth and a third of the compute can still do more work per joule. If energy per token ever becomes the target, the ANE returns as a candidate for short-prompt work that fits the 2 MB window. Section 11 keeps that as lane 5, and it is a measurement this project cannot yet run.
What this project has not measured. Every ANE figure on this page is Bryngelson's, on a part whose GPU roofs are about twice this machine's. No ANE measurement has been reproduced here, and none can be with the current tooling: there is no MLX path to the block, and the third-party drivers that reach it use a private, unsupported API. The ratios do not transfer unchanged, because the two blocks scale differently with the part. Apple lists the same 16-core Neural Engine for every tier of a generation, while the GPU on this base M5 has about half the roofs of Bryngelson's part. If the ANE figures hold on this machine, the inferred gaps here are 5.0× at decode (120.4 against 24.1 GB/s) and 1.6× at prefill (15.4 against 9.6 TFLOP/s): smaller than on the larger part, and still against the ANE at every intensity.
MLX ships two prefill attention kernels: a Steel simdgroup kernel that runs on the general ALUs, and a NAX variant of it that runs on the Neural Accelerators. On M5 the NAX variant is selected whenever metal::is_nax_available() is true, and it instantiates a fixed set of tile shapes. Head dimension 256 did not fit that tile until August 2026. How it was made to fit, and what it still costs, is the open work.
audit/sdpa_string_probe.py, telemetry/sdpa_string_probe.csv (2026-08-27), visible-work metric per the METHODOLOGY rule of that date; the evolved kernel re-validated on the same metric the same day · read from source MLX main 4dc0210, scaled_dot_product_attention.cpp: the NAX kernel is instantiated for hd ∈ {64, 96, 128, 256}; hd=256 runs fused only when qL ≥ 1024 and the mask is causal or an array; hd=80 and hd=192 stay off the NAX path · published ml-explore/mlx #3842 (the head-split kernel), #4416 (the array-mask gate, this project's)One query token per pass. The parallelism has to come from somewhere other than the output, because there is only one row of output.
Measured decode throughput on this machine, as a share of the 120.4 GB/s bandwidth roofline:
| Decode shape | % of bandwidth roofline | Note |
|---|---|---|
| Qwen GQA, 32k KV | 64% | at long context, where attention dominates decode time |
| gpt-oss hd=64, many heads | 49% | short heads, launch overhead a larger share |
| MoE FFN, MLX path | 38–45% | expert gather runs against the 7.0 GB/s random-access floor |
The MoE gather is the largest single structural loss on this chip. Streaming reads sustain 120.4 GB/s; random gathers collapse to 7.0 GB/s, a 17× cliff. Any kernel that turns scattered expert reads into contiguous streams of a megabyte or more is recovering that factor directly, and that is worth more than any arithmetic change available in decode.
This table is read from MLX main (4dc0210, 2026-09-02); the July rows were confirmed by building v0.32.0 with an instrumented dispatch. The NAX kernels are not a separate family from Steel: the headers are steel/gemm/nax.h and steel/attn/nax.h, described in-source as "NAX Steel with new tiles". Same Steel structure, different tiles, different datapath.
| Condition | Family selected | Runs on | Ceiling |
|---|---|---|---|
is_nax_available() and hd ∈ {64, 96, 128} | sdpa_full_self_attention_nax | Neural Accelerators | 15.4 TFLOPS |
| hd = 80 excluded by name | Steel simdgroup path (non-NAX tiles) | general ALUs | 6.1 TFLOPS |
| hd = 256, qL ≥ 1024, causal or array mask #3842, #4416 | attention_nax_dsplit | Neural Accelerators | 15.4, reaches 57.5–61.5% |
| hd = 256 otherwise; hd = 192 qL > 8 | two mx.matmul + softmax | Neural Accelerators | 15.4, reached 34–44% at hd=256 July metric |
| architecture gen < 17 M4 and earlier | Steel simdgroup path | general ALUs | n/a on this part |
| decode, any hd | mx.fast.sdpa decode path | bandwidth-bound | 120.4 GB/s |
The counterintuitive row is the unfused one. The "unfused fallback" at hd=256 is not a slow path. It is two large matmuls that run on the Neural Accelerators. A fused simdgroup kernel that replaces it measured 0.37×: correct, and 2.7× slower, because it trades NA silicon for ALU silicon. Fusion is not automatically a win when the operands sit on different datapaths.
The items the census left open, and the ones it closed.
| Open item | Status | Why it matters |
|---|---|---|
| packed int4 sustained rate | not measured | first-class operand type in the tensor-op API; int8 measures 18.6 TOPS, packed int4 has no measured rate yet |
| SLC effective size | measured | 4–8 MiB at ~80 ns from the chase ladder; capacity, not TLB, per the page-local control |
| GPU clock under sustained NA load | not measured | needed to convert the 15.4 TFLOPS aggregate into a per-cycle figure for this part |
| fp8 | ruled out | absent from the toolchain (build 17F109) at the API level |
| hd=256 prefill ceiling | open | the upstream fused kernel measures about 60% of ceiling; incremental fixes to it cap at 71.5% / 64.5% (4k / 32k); no kernel structure has measured 80% on the visible-work metric, so about 70% is the evidence-supported ceiling today |
| ANE as an energy target | open | throughput is settled against the ANE by Bryngelson's roofs (section 06: 9.5× slower at decode and 3.2× at prefill on Bryngelson's part, about 5× and 1.6× inferred for this base M5). What stays open is energy, where the same study measures it 13× more efficient, and only for work inside its 2 MB operand window. No MLX path, no bf16, and an unsupported private API block that measurement here |
Ranked by expected value against this machine's measured limits, not by how interesting they are. Decode is bandwidth-bound and prefill is compute-bound, so a lane only matters if it moves bytes on the decode side or intensity on the prefill side.
| # | Lane | Why it could pay | Cost to find out | Models that will benefit |
|---|---|---|---|---|
| 1 | Contiguous expert gather for MoE | Random gather measures 7.0 GB/s against 120.4 GB/s streaming: a 17× cliff, the largest structural loss on the chip. Decode is bandwidth-bound, so this sits directly on the critical path. MoE decode currently runs at 38–45% of the bandwidth roofline. | kernel work; the target shape is already known | Mixture-of-experts models at decode, at every context length: Qwen3-Coder-30B-A3B, Qwen3.6-35B-A3B, gpt-oss-20b. The gain is tokens per second during generation; their decode runs at 38–45% of the bandwidth roofline today. |
| 2 | Audit MLX's GEMM tile sizes against the intensity bound | Section 04 derives AI = BM·BN/(BM+BN), so a square tile must be 256×256 to reach the 128 FLOP/byte ridge. A 128×128 tile caps out at AI 64, which is 50% of peak by tiling alone. If MLX ships smaller tiles than the bound allows, prefill GEMMs lose throughput for a reason that is arithmetic, not engineering. |
a source read plus one sweep | Dense models at prefill, so time to first token: Llama 3.1 8B, Qwen3 14B, Mistral 7B, Gemma 3 12B. The share is largest on short and medium prompts, where the matrix multiplies are most of prefill; by 32k tokens attention takes about two thirds of it (68%, measured on Qwen3-Coder-30B). |
| 3 | Measure the int8 and packed-int4 operand rates | Measured int8 is 18.6 TOPS against 15.4 TFLOPS fp16, a 1.21× gain for half the operand bytes. Packed int4 is a first-class operand type in the tensor-op API and has no measured rate at all. Decode is bandwidth-bound, so an operand format that halves or quarters weight bytes acts directly on the binding constraint. | one probe revision | Any model at prefill, if int8 (18.6 TOPS measured) or packed int4 beats fp16 on the Neural Accelerator path. At decode only where a native int4 operand saves the in-kernel unpacking the current 4-bit kernels do, which on this 24 GB machine means every model served. |
| 4 | hd=256 prefill: close the remaining gap on the NA path | Two steps are done. MLX #3842 (merged 2026-08-19) runs hd=256 causal prefill fused on the Neural Accelerators by splitting the head dimension across two simdgroups. This project's #4416 (merged 2026-09-02) extends that path to array masks, the form sliding-window layers use: 1.44–1.56× at kernel level, 1.05–1.06× time to first token on Gemma-3-12B. What remains: the fused kernel sits at 61.5% of ceiling at 4k and 57.5% at 32k against about 79% at hd=128, an accumulator-pressure cost. Incremental fixes measure a ceiling of 71.5% / 64.5%; no structure has measured 80% on the visible-work metric. 60 → 70% prices at about 1.05× time to first token at 32k and 1.08× at 128k on Qwen3.6-35B. | a kernel restructure (campaign-class); the incremental route is measured and capped | Models with head dimension 256 at prefill, so time to first token, on long prompts: the Qwen3.5 / 3.6 / 3.8 full-attention layers (Qwen3.6-35B: 10 of 40 layers; attention is 19% of its prefill at 32k and 48% at 128k) and Gemma 3 4B and 12B (one global layer in six; their sliding-window layers already gained from #4416). |
| 5 | Locate the ANE crossover for this project's shapes | Section 06 settles throughput against the ANE: 9.5× slower at decode and 3.2× slower at prefill on Bryngelson's part, about 5× and 1.6× on this base M5 if the block is the same, with its speed advantage confined to a 2 MB operand working set. This is not a throughput lane. What survives is ~13× energy efficiency. The measurement worth doing is where the energy crossover falls for these sequence lengths, and whether energy is a target at all. | one measurement, blocked on tooling | None, for throughput, at these prompt lengths. If energy per token becomes a target: short-prompt work on any model, where the operand working set stays under 2 MB. |
| 6 | Separate cache residency from TLB reach · DONE | The gather window is capacity, not translation reach. Translation costs appear only from about 64 MiB, so lane 1 should size its contiguous blocks against the 4–8 MiB SLC window. | done: the last probe in section 02, 2026-09-01 | The same mixture-of-experts models as lane 1; this result sets the block size lane 1 should use. |
audit/value_model.py on Qwen3-Coder-30B; the 38–45% decode figure is section 08Start with lanes 1 and 2. Lane 1 moves bytes on the bandwidth-bound side, which is decode. Lane 2 raises arithmetic intensity on the compute-bound side, which is prefill. Everything else on this list is either already running, a low-cost check, or a measurement that narrows a question rather than moving a number.