Flint Inference on Apple Silicon: rec Megakernel and Tempo9 Compared

2026-10-27

English · 中文版

rec, released by Recursive alongside Flint, is an Apple Silicon inference engine specialized for this model. Its persistent Metal megakernel performs the model's decoding work within a single dispatch. For single-stream decoding, that approach is remarkably effective.

We tested rec alongside Tempo9, a general-purpose inference engine, on the same M5 Pro MacBook Pro. We measured single-stream decoding, concurrent serving, cold prefill, cross-session prefix-cache reuse, and CPU use, then examined the underlying execution with per-layer profiling and controlled Metal experiments.

Rather than ranking the engines, we focus on two questions: what makes rec's single-stream decoding so fast, and what happens when real workloads switch between conversations? The benchmark setup and methodology are provided near the end.

1. Performance: Different Strengths at Different Workloads

All primary benchmarks were run on an M5 Pro MacBook Pro with 24 GB of unified memory.

To make the comparison meaningful, we used model builds with nearly identical weight traffic per generated token: approximately 1.385 GB for Tempo9 and 1.39 GB for rec.

This matters because differences in quantization can easily overshadow differences in inference-engine efficiency.

Single-stream decoding: rec is 26% faster

For a single request generating 256 tokens with greedy decoding, rec achieved 151–155 tok/s, compared with 123–127 tok/s for Tempo9.

That makes rec approximately 1.26× faster in single-stream decoding.

This is an impressive result. For workloads dominated by one local conversation, rec's highly specialized execution model is effective. But single-stream throughput is only part of the picture.

Concurrent inference: Tempo9 pulls ahead

rec is designed around a single active session. It does not support continuous batching of independent requests; additional requests must wait for the active request to complete.

Tempo9 uses continuous batching, dynamically combining requests at each decoding step so that multiple sessions can share the GPU.

We tested 1, 2, 4, 8, and 16 concurrent clients, each requesting 256 generated tokens.

With one client, rec leads. At two clients, Tempo9 moves ahead in aggregate throughput. The gap then widens as concurrency increases.

At 16 clients, Tempo9 reaches 378 tok/s, compared with 141 tok/s for rec — approximately 2.68× the aggregate throughput.

The slowest request also completes substantially earlier: 12.2 seconds with Tempo9 versus 29.0 seconds with rec.

These results illustrate a fundamental difference in design priorities. For a dedicated, single-user inference session, maximizing individual token throughput may be the primary objective.

But a local AI system increasingly runs more than one workload: an IDE assistant, an autonomous agent, background tool calls, and interactive conversations may all compete for GPU time. In those scenarios, efficiently scheduling multiple requests can matter as much as the speed of any individual request.

Prefill: it depends on the length

We compared cold prefill latency — no cached prefix — at four prompt lengths.

At 562 tokens, Tempo9 completes prefill in 0.37 seconds, versus 0.76 seconds for rec — about 2.1× faster. At 1,971 tokens, rec has a small advantage: 0.80 versus 0.84 seconds.

At 8,022 tokens, rec is ahead: 2.8–3.0 seconds against 3.7–4.1 for Tempo9 over three interleaved rounds. (An earlier single run had rec at 4.12 seconds; it was taken under memory pressure and has been replaced.) At 9,289 tokens the order flips again: Tempo9 4.5–4.9 seconds, rec 5.3–5.6. rec prefills in 4,096-token chunks, and its cost rises steeply once a prompt needs a third chunk; Tempo9's cost grows roughly linearly with length.

Neither engine leads across the board. Where prefill matters most for an agent is not the cold case but the repeated one, which the next section measures.

Prefix caching: consecutive turns versus interleaved conversations

Prefix caching avoids recomputing previously processed context. However, continuing one conversation and switching among independent conversations are distinct workloads and should be evaluated separately.

rec's PromptCache retains the current conversation's tokens and KV state, matching prefixes token by token. Tempo9 uses shared prefix caching in 128-token spans. We first measure a continued conversation, then test whether the cache survives a switch to another conversation.

Test 1: One continuing conversation — rec achieves an excellent cache hit

We first send a document of roughly 8,000 tokens to conversation A, then append another question within the same conversation. This resembles a multi-turn chat or an agent appending tool results.

rec's single-session prefix cache performs extremely well. It reuses all 8,022 existing tokens, prefilling only the 40 newly appended tokens. Tempo9 caches complete 128-token spans, so the trailing partial span is recomputed. In this test, rec achieves both a more complete hit and lower prefill latency.

rec also uses prepare_prefix() to include generated tokens in the cache for the next turn. It can roll back to a checkpoint when a prompt branches from an earlier point; that is a separate behavior from switching between independent conversations.

Test 2: Interleaved conversations — does a previous prefix survive?

Now consider two independent conversations sharing one server. We issue requests in A → A → B → A order: start A, continue A, start B with a different document, then return to A. Each sequence was tested in three rounds; the chart uses engine-reported prefill times.

The two charts separate continued-session cache hits from cross-session retention. The second chart focuses on returning to A after serving B.

When we return to A, rec's single PromptCache has been replaced by B, so it must recompute A's context: 2.86–2.88 seconds, comparable to cold prefill. Tempo9 still reuses A's 7,936 cached tokens, taking 0.28–0.31 seconds.

This does not mean rec lacks prefix caching. It handles a single growing conversation exceptionally well; the limitation of the tested implementation is retaining prefixes from multiple independent conversations at once. That matters when a main agent and subagents alternate, or several chat windows share a server.

In the third test round, rec's final request also encountered a Metal out-of-memory error; the corresponding 9.3k-token prefill peaked at 23.2 GB on this 24 GB machine.

CPU usage during decoding

The engines also differ in how much CPU time they consume while the GPU is decoding.

During decoding, Tempo9 consumes approximately 0.15–0.17 CPU cores, compared with 0.87–1.00 CPU cores for rec.

Sampling shows that much of rec's CPU activity comes from an MLX stream thread polling inside Fence::wait while waiting for GPU completion. This is with rec's default --fast-sync (MLX_METAL_FAST_SYNCH=1), which trades CPU time for lower CPU–GPU synchronisation latency; we did not change it.

Tempo9 takes a different approach. Its inference engine is implemented in C++. Generated tokens move to the serving layer through a lock-free queue, while GPU completion is coordinated through command buffers rather than a busy-waiting loop.

The CPU work required to encode the next decoding step — approximately 1.3–1.5 ms per token — overlaps with GPU execution of the current step.

This is particularly relevant for local AI agents. The CPU is not dedicated to inference: it also executes tools, manipulates files, runs browser operations, and supports the rest of the application. Lower inference-related CPU overhead leaves more resources available for those tasks.

2. Where Does rec's Single-Stream Advantage Come From?

To understand rec's lead, we examined its open-source implementation and profiled its execution at the layer level.

A conventional GPU inference engine executes a Transformer through a sequence of operators: RMSNorm, QKV projection, RoPE, attention, output projection, MoE routing, and expert computation.

Each operator is typically implemented by one or more GPU kernel dispatches, with the execution graph determining their ordering.

For readers familiar with CUDA: A Metal SIMD group roughly corresponds to a CUDA warp, a threadgroup to a thread block, and threadgroup memory to shared memory. Metal's threadgroup_barrier() is analogous to CUDA's __syncthreads(): both synchronize only within a group. A Metal compute dispatch can be thought of as a CUDA kernel launch.

On the M5 Pro, rec launches 36 threadgroups of 256 threads each—roughly 36 CUDA-style blocks with 256 threads per block. To synchronize across threadgroups, rec uses device-scope atomics and spinning, which requires participating groups to be simultaneously resident. That is not a general grid-wide synchronization guarantee. Apple GPU cores are not exact equivalents of CUDA SMs, nor are Metal command buffers interchangeable with CUDA streams.

The diagram below shows the M5 Pro execution hierarchy, from unified memory to SIMD groups and threadgroups. The synchronization boundaries are central to rec's design.

In Metal, a compute dispatch comprises threadgroups, each made up of SIMD groups. Threads within a threadgroup can share on-chip memory and synchronize with threadgroup_barrier(); separate threadgroups cannot use that primitive to synchronize with one another. rec instead coordinates them through device-scope atomics and spinning barriers, an approach that depends on the participating groups being resident together.

For the underlying model, see Apple's threadgroup documentation, the Metal Shading Language Specification, and WWDC's Apple GPU compute session.

rec takes a different approach. On M5 Pro, it launches 36 persistent threadgroups, each containing 256 threads. These threadgroups collectively execute Flint's 49 Transformer layers inside a single Metal dispatch.

Synchronization between threadgroups is implemented using device-scope atomic counters and spinning barriers.

The following diagram summarizes a decode step within one dispatch, showing where resident threadgroups synchronize and where operations are fused.

The more consequential optimization happens within each threadgroup.

A threadgroup responsible for an attention head can perform projection, normalization, and KV writes while retaining intermediate values within the same execution context.

This allows rec to reuse data in registers and threadgroup memory instead of repeatedly moving intermediate results through separate kernels.

Every operator of the layer lives inside that one dispatch, not only the attention path: the input RMSNorm, the post-attention RMSNorm, the MoE router, the top-6 selection, the expert gate/up and down projections (layers 0 and 1 have a dense MLP instead, and the kernel carries a branch for it), and after the last layer the final norm, lm_head and argmax. The diagram above shows the two RMSNorms and the top-6 without a barrier of their own, which is deliberate: each threadgroup holds the full 2048-wide residual in its registers (eight values per thread) and recomputes the norm itself — a few thousand multiply-adds per threadgroup in exchange for not waiting on the other 35. The experts run as teams, and the hand-off between gate/up and down is team-local rather than a grid-wide barrier. The barriers that remain are the ones where data genuinely crosses threadgroups: after qkv, after the attention partials, after the merge, after o_proj, after the router and after the experts.

The architecture is tightly coupled to Flint's execution structure, which allows optimizations that would be harder to express in a fully general computation graph.

Per-layer profiling: the gap is in attention

To understand rec's single-stream advantage, we split each Transformer layer into attention and MoE stages and measured their execution times separately.

In attention, rec takes approximately 44 µs per layer, versus 63 µs for Tempo9. The MoE stage shows the opposite pattern: Tempo9 takes 94–100 µs, compared with approximately 109 µs for rec.

rec's advantage is concentrated in attention, not expert computation. Its megakernel closely integrates normalization, RoPE, KV writes, and adjacent operations, reusing intermediate data within threadgroups and reducing transfers and synchronization between separate kernels.

Across Flint's 49 layers, small differences accumulate into meaningful token latency. The per-stage rec measurements come from a profileable kernel, while the 151–155 tok/s result at the start of this post comes from its further optimized path; they should not be treated as directly additive. Next, we use a controlled experiment to isolate the value of removing dispatch boundaries alone.

3. The Real Benefit of a Megakernel Is Fusion, Not Simply Fewer Dispatches

At first glance, the explanation for rec's performance seems straightforward: one GPU dispatch should be cheaper than many. But how much of the difference actually comes from dispatch overhead?

To answer that question, we built a controlled Metal experiment using the same device functions for one attention half-layer.

The first path uses seven separate dispatches within a single command buffer. The second executes the same computation inside a persistent kernel, using atomic barriers between stages.

The two main paths had nearly identical latency.

We also measured synchronization costs. On our M5 Pro, a cross-threadgroup barrier takes approximately 1.0–1.5 µs, while a kernel dispatch boundary costs about 1.4 µs.

Replacing multiple dispatches with one persistent dispatch does not automatically make the computation faster.

The improvement depends on what can be fused within each stage. rec benefits from keeping intermediate values close to the computation, eliminating redundant memory operations, and combining work that would otherwise require separate kernels.

Those are valuable techniques even without adopting a full-model megakernel.

This suggests that selective operator fusion can capture some of the same benefits without adopting a full-model megakernel.

4. Why Tempo9 Does Not Use a Full-Model Megakernel

A megakernel is a natural fit when the model architecture, execution shape, and target hardware are known in advance.

Tempo9 has a different goal. It is designed as a general-purpose inference engine supporting multiple model families, quantization formats, hardware backends, and concurrent workloads.

Tempo9 targets more than Apple Silicon. It also supports NVIDIA CUDA GPUs and ARM CPU-only devices, while sharing a model execution framework across architectures including Qwen, Gemma, Llama, Mistral, and Flint. Its engine core is implemented in native C++. On the CUDA backend, CUDA Graphs allow the inference execution path to be captured and replayed where applicable, reducing host scheduling and repeated launch overhead. This is a CUDA-specific execution mechanism; it does not imply that Metal or ARM CPU paths use the same GPU graph API.

This is a tradeoff between runtime performance and software complexity. A specialized megakernel can fuse work more aggressively for one model and target device. A reusable C++ execution framework, shared operators, and a CUDA Graph execution path make it easier to support different model architectures and hardware without maintaining a separate full-model megakernel for each combination. That flexibility is not free, but it is an intentional engineering choice.

The following diagram traces a decoding step through Tempo9's scheduler and operator chain. Unlike rec's fixed work allocation, the active batch can change between steps.

These operator boundaries help in three areas.

Dynamic batching

Real-world agent workloads are inherently dynamic. Requests enter and leave the active batch. Some are performing prefill, others are decoding, while others are waiting for tool execution.

Tempo9 rebuilds the active batch at every decoding step and selects the appropriate execution kernels based on the current shape. Small batches can use GEMV-oriented execution, while larger batches can switch to tiled or grouped GEMM paths.

rec's megakernel is specialized for a fixed execution structure. Its lack of continuous batching is a direct consequence of that design.

Shared kernels across model architectures

Tempo9 uses a common collection of computational kernels across model families, including Qwen, Gemma, Llama, Mistral, and Flint.

The model architecture defines the operator graph, while reusable kernels implement its computational building blocks. This makes it possible to support new architectures without writing an entirely new GPU inference kernel for each model.

A specialized megakernel offers more opportunities for model-specific fusion, but it also increases the amount of architecture-specific and hardware-specific implementation work.

Features that change the execution graph

Prefix caching, speculative decoding, KV management, and dynamic scheduling can all change how inference work is organized.

Tempo9 needs these features to coexist with the same underlying execution system. Maintaining operator boundaries provides flexibility without generating a separate megakernel for every execution variant.

This flexibility has a measurable cost in batch-one decoding, but it also enables the concurrency, model coverage, and serving characteristics shown in our benchmarks.

The two approaches are not mutually exclusive: a general-purpose engine can use local fusion without requiring a separate megakernel for every model.

5. Benchmark Setup and Methodology

Hardware and software

• Hardware: MacBook Pro with M5 Pro, 24 GB unified memory
• Operating system: macOS 26
• Tempo9: version 1.2.0, engine commit 754721a8f
• rec: upstream main at 0dcf43c (October 22, 2026) plus two small local changes, published as thinkspread/rec, branch local/flint-e3-24gb. 4eacf0a removes the 48 GiB startup check, accepts the INT3-expert checkpoint and wires the weights with mx.set_wired_limit; without it rec does not start on a 24 GB Mac. 325d244 raises the server's admission semaphore from 9 to 32 so that 16 clients can connect for the concurrency test. Neither change touches the kernel, the scheduler or the prefix cache.
• Measurement period: October 23–26, 2026

Quantization and weight traffic

The primary speed comparison uses approximately equal bytes read per token, rather than similarly named quantization formats.

The published Tempo9 build uses a different quantization configuration and should not be confused with the equal-byte benchmark build. Model quality was evaluated separately against a Q8_0 reference.

Workload

Single-stream decoding used greedy generation of 256 tokens from short prompts, ending at approximately 300 tokens of context. Concurrent tests used distinct prompts, with each client requesting 256 generated tokens.

Measurement procedure

All comparisons were performed on the same machine. We stopped the local CI runner during testing and alternated benchmark configurations within the same session for at least three rounds.

Single-stream throughput was taken from the engines' counters; concurrent throughput was measured at the client. We monitored memory state, since compressed-memory pressure on macOS can substantially affect absolute performance.

These controls do not eliminate every source of variation, but make comparisons more informative than isolated runs.

6. Try Flint with Tempo9

Install Tempo9:

brew install thinkspread/tap/tempo9

Or use the installation script:

curl -fsSL https://raw.githubusercontent.com/thinkspread/tempo9/main/install.sh | bash

Start Flint:

tempo9 --hf thinkspread/Flint-Q6K-IQ3XXS-GGUF

The server listens on localhost:11435 and provides OpenAI- and Anthropic-compatible APIs.

You can also load a local GGUF file:

tempo9 --gguf <file>

7. Conclusions: A Single Megakernel Is a Promising Starting Point

rec demonstrates the potential of deep, model-specific optimization on Apple Silicon. Its persistent Metal megakernel delivered 151–155 tok/s in our M5 Pro single-stream test, approximately 1.26× Tempo9's throughput at comparable weight traffic per token. Per-layer profiling also shows that careful fusion and data reuse within the attention path can produce meaningful gains.

Executing a model's decoding path as a single Metal kernel is a promising starting point. It is a useful execution model to explore and a clear demonstration of what specialization can achieve. At the same time, one dispatch is not the whole explanation: in our experiments, cross-threadgroup barriers and kernel dispatch boundaries had comparable costs. The more important question is what computation and data movement fusion eliminates.

Concurrent serving and prefix-cache management deserve equally serious consideration as an inference engine develops into a system for real workloads. The tested version of rec does not support continuous batching, so aggregate throughput under multiple clients does not scale like a dynamically batched engine. rec's prefix cache is particularly effective for a single, linearly growing conversation: on A's second turn it reused all 8,022 tokens and needed to prefill only 40 new tokens, taking 0.13–0.14 seconds versus Tempo9's 0.22 seconds. The limitation is cross-session retention: in the A→B→A sequence, the tested rec implementation replaced A's PromptCache with B's, requiring a full prefill on returning to A. That is a limitation of the current single-PromptCache implementation, not an inherent restriction of the megakernel approach.

For a dedicated Flint conversation where sustained decode speed is the priority, rec's design is compelling. For multiple agents, interactive sessions, and workflows that alternate between inference and tool calls, scheduling and cross-session KV-cache reuse also have a substantial impact on end-to-end experience.

These directions need not be mutually exclusive. rec has taken a valuable step toward efficient single-stream GPU execution; extending that foundation to concurrent requests and reusable prefixes across sessions is an equally interesting area for further work.

We thank Recursive for releasing rec under the Apache-2.0 license. Access to the implementation makes it possible to examine these architectural tradeoffs on technical grounds rather than relying on throughput numbers alone.