tile-rs on Apple GPU: our Metal backend against our Mojo backend

28 August 2026

TL;DR — tile-rs is a Rust project for writing accelerator kernels that are type-, concurrency- and resource-safe by construction, and lowering them to many backends. Two of those target Apple GPU: one emits Metal Shading Language, one emits Mojo. On a Mac Studio (M1 Ultra, 128 GB), across sixteen kernels — five primitives and eleven from the decode path of DeepSeek-R1-Distill-Qwen-1.5B — our MSL backend is ahead on all sixteen, from 1.18× to 2.99×.

That result is specific to that machine. On an M4 Mac mini the same sixteen kernels put MSL ahead on nine, level on six, and behind on o_proj at 0.68×. The kernels are byte-identical and every one passes the same value check on both; what changes is the machine.

Both columns are ours, so this compares two tile-rs backends — it is not a measurement of Mojo. Adding a third arm makes that concrete: MAX's own Apple GEMM, which accumulates through simdgroup_matrix hardware MMA, beat both of ours at every matmul size above 256³ — by 3.1× over our corpus kernel at 1024³ on the M1 Ultra, and 3.9× on the M4. Reading their kernel showed us why: rescheduling our own mma lowering the same way now puts it ahead of both, at every size we measured on both machines.

Why this comparison is possible now

On 11 August 2026 Modular shipped Mojo 1.0 with stability guarantees, and on 18 August they open-sourced the compiler and toolchain under Apache 2.

That second step is what makes this post possible. Before it, Mojo's numbers were Modular's to publish. Now anyone can build it, run it, and check.

tile-rs is a much smaller thing, open-sourced in late June 2026. It is a Rust dialect for accelerator kernels whose defining property is that the compiler refuses to emit a kernel that could race, overflow a buffer, or read a tile it does not own.

Both projects target Apple GPUs. That overlap is the whole experiment.

The head-to-head

Sixteen kernels on a Mac Studio (M1 Ultra, 128 GB), at the project's strict shapes — the same widths every other table here uses, so a kernel's row can be compared across backends.

The protocol is paired. Each round runs A/B/B/A — MSL, Mojo, Mojo, MSL — so drift in machine state falls on both sides equally; five rounds; the statistic is the median of per-round ratios, not a ratio of medians, so a slow round cannot be charged to one arm. Both sides go through the same harness, the same shapes and the same value check, run before anything is timed.

It is wall-clock on both sides, deliberately. Metal can report gpuStartTime/gpuEndTime; Mojo on this backend cannot — DeviceEvent constructs and then fails at link with Symbols not found: _AsyncRT_DeviceContext_enqueue_event. Timing one side on the device clock and the other on the wall would compare two different instruments, so both are timed the same way.

The same kernels on other Macs

Run on a MacBook Pro (M2 Max, 64 GB) and a Mac mini (M4, 16 GB), same script, same kernels, five rounds:

Correctness is identical everywhere — 16/16 pass on all three, at the same relative error per kernel — so this is about the machine, not the code.

The M4 is the interesting one. gate_up_silu takes 983 µs there against 143 on the Studio: it moves 110 MB of weights per launch and a 16 GB Mac mini has nowhere near the bandwidth to hide that. When a kernel is pinned to memory both backends hit the same wall and the code generator stops mattering, which is what the six level rows are. o_proj is different: at around 75 µs against Mojo's 51 it measures behind consistently, and we have not found the cause. Threadgroup size from 64 to 1024, three shapes of cross-simdgroup reduction, and compiling one code path instead of two all leave the median where it was; batching every launch into a single command buffer, which is what the Mojo harness effectively does, is worth about 5%. The M4 is also the noisiest of the three machines — the same kernel measured 55 µs and 70 µs in two sessions an hour apart — so we would not read a precise ratio off that row.

Read the fourth column before the third. At these shapes most of this corpus is launch-bound on both sides: Metal's own device clock puts vec_add at about 8 µs of GPU time against ~31 µs on the wall, and Mojo's launch floor is around 30 µs for the same kernel. So the 1.3× band is a real measurement of two launch paths and says almost nothing about the generated code.

The rows worth arguing about are the ones with work in them — down_proj, argmax, attn_gqa, gate_up_silu — and they range from 1.18× to 2.99×. matmul is not among them: at 64³ it is 0.5 MFLOP, so that row measures dispatch too. The next section takes it to sizes where the arithmetic shows. gate_up_silu at 1.18× is the honest counterweight to argmax at 2.99×.

Method note. On a shared machine some rows come back with a spread larger than the effect. The harness refuses those — it prints INCONCLUSIVE rather than a number when the effect is under the spread — and the table above is assembled from the cleaner of two runs per kernel, selected on spread and never on ratio. Each kernel is its own paired measurement, so that is choosing the more precise instrument, not the more flattering answer. Every row here sits between 2% and 24% spread.

A third kernel, and what reading it was worth

Both matmul kernels above have the same shape: one threadgroup per output row, each thread owning a strip of columns, a scalar loop over K. Neither emits Apple's simdgroup_matrix instructions — though our Metal backend can: mma is a typed operator in the calculus, and its Metal lowering to simdgroup_multiply_accumulate has been there all along. The corpus kernel simply does not route through it.

MAX's kernel library schedules that instruction properly, so we ran its Apple GEMM through the same harness as a third arm. It beat both of ours — by 3.1× at 1024³ on the M1 Ultra and 3.9× on the M4. The gap was a scheduling gap, not a missing primitive, which made it ours to close.

The first attempt closed most of it and stopped: staging a 128×64 block through threadgroup memory fixed the redundant reads but paid a barrier pair on every K step, which at K=1024 is 128 barriers a block. It landed at 1.16× and 1.32× behind — better, and still behind.

What actually closed it came from reading their kernel, which is what open sourcing it is for. Their design is the opposite of staging: no threadgroup memory and no barriers at all, fragments loaded straight from device memory into registers, with an unguarded fast path for subtiles that lie wholly inside the matrix. Metal's simdgroup_load reads the device address space directly, so the same structure expresses in MSL without borrowing their per-lane fragment layout. Staging stays only where it buys safety — blocks straddling an edge, and the K remainder, take a guarded path that zero-fills anything out of range:

At 1024³ that is 4.50× our own scalar path and 1.53× MAX's kernel; on an M4 mini it is 2833.2 / 686.9 / 725.5, so 4.12× and 1.06×. The lowering is now ahead at every size we measured on both machines, including 128, where the staged version had lost to the scalar path.

Rescheduling also exposed a defect worth naming. The original lowering accumulated in simdgroup_half8x8, and it was faster than the corrected version — because at K=512 it returned relative error 1.7e-2 against a host reference and failed the same value check every other row in this post passes. Checks at small K had let it through for a long time. It accumulates in f32 now, at 4.8e-4.

Three things follow, and they matter more than the first table.

Our emitted Mojo is not a measurement of Mojo. It is naive code that happens to be written in Mojo, and a library kernel is the fair opponent. Where the two differ by 9×, the difference is the emitter, not the language.

A ratio is only as good as the shape it was taken at. The corpus runs at decode shapes, where most kernels are launch-bound and the code generator barely registers. That is worth saying plainly rather than leaving in a column header.

A number that loses is worth more than one that wins. Being 3.9× behind was what sent us to read the other implementation, and the lever was in its first twenty lines. None of the four hypotheses we tested before reading it were right.

Method, for this table. The M1 Ultra was under other load throughout. The measurement is paired and its scalar arm reproduces an independent paired run to within 0.15% at every size, which is the reason the rest of the column is quoted at all; unpaired single-shot readings on the same machine varied by a factor of four and are used nowhere here.

Holding the engine fixed

The table above compares two of our own backends, so it cannot say whether our kernels are any good — both sides are ours. For that, hold an engine fixed and swap only the kernel.

So we took llama.cpp, built it from source, and dropped our emitted matvec into ggml-metal behind an environment switch. Same binary, same scheduler, same KV cache layout, same sampling, same dispatch geometry. The kernel is the only thing that differs.

0.951 — 4.9% slower, conclusive over six ABBA rounds with byte-identical greedy output. It is the most useful number here precisely because it is the one we would not have chosen.

Then the mirror, because one direction proves nothing about a kernel and everything about an engine. We lifted ggml's kernel_mul_mv_f16_f32_4 verbatim — its body, its reduction helper, its function constants, its args struct — into our engine, and matched its launch geometry exactly:

0.955 the other way. Each kernel is about 5% faster inside the engine it was written for, in both directions, with the sign test unanimous both times.

A kernel and its host are co-designed: moving either across costs roughly the same. Run singly, the first table reads as a deficit and the second as a win, which is why both have to run.

How? Three claims

One: the kernels are safe, in three specific senses.

• Type safety. No raw pointers, and nothing that needs a runtime check to stay true.
• Concurrency safety. No data races, with no synchronisation the programmer has to remember to add.
• Resource safety. A kernel cannot allocate more than the buffer it was given. Not theoretical: on other targets in this project, backends that checked a budget they had merely written down — rather than one the type system enforced — got a hardware capacity limit wrong, and the miss surfaced anywhere from the assembler to the device mid-run.

Two: hardware affinity is not optional, and not uniform.

GPUs, TPUs and NPUs have genuinely different architectures, and no single emitted implementation suits all of them. What they do share is enough structure to be captured by one type system and one set of invariants. tile-rs generates a different implementation per target, keeping the operator's meaning while satisfying that target's memory, concurrency, synchronisation and capacity constraints.

Three: the two directions compose.

Lowering goes from types and invariants to a hardware-affine implementation; it must preserve the safety properties while introducing hardware detail. Lifting goes the other way — from an optimised, target-specific kernel back to types and invariants — and it must re-establish safety. The payoff is what happens when you lift and then lower again to a different target: an optimisation discovered on one backend transfers to the others.

In practice this has been a magpie exercise: tile-rs borrows from the tile abstractions its neighbours have built, Mojo's included, then lifts and lowers repeatedly. Adding safety to those ideas has, repeatedly, also made them faster.

Credit where it is due

Mojo's abstraction level holds up well on safety, and here is the evidence for that from outside the timing tables. Every kernel in this corpus exists in two readings: the shapes as they were first captured, some of them degenerate and ill-formed, and the same kernels after those shapes were fixed. Our other backends emit different code for the two readings — they were sensitive to a malformation they had no business being sensitive to. The Mojo operator is byte-identical across both. It never had to care.

That is a real design win, and hats off to the Modular team for it.

The comparison we have now partly run, and an open offer

Modular's position is that MAX, not the Mojo language surface, is the layer for AI/GPU work, and that is the right distinction: our emitter targets the Mojo language because that is what an emitter can target, and a language-level kernel comparison says nothing about the performance a MAX user would get.

For matmul we have now done this ourselves, with MAX's own Apple GEMM, and it beat both of our backends — the section above. That is one kernel of sixteen, and the library's other Apple paths are gated to M5 hardware we do not have, so the rest of the corpus still has no MAX arm.

The offer is genuine for those. Replace our emitted Mojo — with hand-written Mojo, MAX graph ops, whatever represents the supported path — and we will run it through the same harness, on the same shapes, on the same machine, and publish the result whichever way it falls, as we just did for matmul. The harness is one script and the kernels are drop-in; see Reproducing.

The kernel corpus and the benchmark entries are published at pu-rs.org. If a kernel here can be beaten — by a better emitter, a better Mojo, or a shape we got wrong — send a pull request to the repository below and the numbers move.

Reproducing

The Apple comparison is a single script. It ships all sixteen kernels for both backends, checks each one against a host reference before timing anything, and prints the table above with the spread as well as the median:

git clone https://github.com/yijunyu/tile-rs-metal
cd tile-rs-metal/benchmarks/apple-mojo-vs-msl
./run.sh

It needs a Mac with an Apple GPU, swift, and Mojo 1.0+ — the kernels ship pre-generated, so running the benchmark does not need a Rust toolchain. (Regenerating them from tile-rs does: Rust 1.85 or newer, the crates being edition 2024, and these were generated with 1.96.) The script reports the machine, its memory, the OS and both toolchain versions alongside the numbers, so a result is self-describing. Dropping different Mojo kernels into kernels/mojo/ is all that is needed to run the comparison against something other than our emitter's output.

Headline figures: Mac Studio, Apple M1 Ultra, 128 GB, macOS 26.5. The cross-machine table adds a MacBook Pro (M2 Max, 64 GB, macOS 26.6) and a Mac mini (M4, 16 GB, macOS 26.5), all three on Mojo 1.0.0 (ed45d567). The sixteen-kernel table is at strict shapes, paired A/B/B/A wall-clock, five rounds, median of per-round ratios. The fixed-engine numbers are llama.cpp b1-ca3d5a3 built from source, DeepSeek-R1-Distill-Qwen-1.5B f16, tg128 and 256-token decode respectively, ABBA-interleaved six rounds with identical generated tokens on both arms. Every kernel in every table is verified against a host reference in the run that timed it.