Apple Silicon is among the most widely used platforms for local LLM inference, yet nearly all inference optimizations are designed and validated on server-class NVIDIA GPUs under batched serving workloads. Which optimizations retain their benefits when moved to Apple Silicon, and what determines their applicability? We evaluate four optimization families: KV-cache pruning, speculative decoding, ported GPU kernels, and post-training quantization. Experiments span four Apple machines under a common single-request decode workload, with matched RTX 4090 comparisons where possible. KV-cache pruning can reach up to 1.75× end-to-end throughput on Apple but can reduce throughput on desktop NVIDIA GPUs. Under the common benchmark, all evaluated speculative decoding methods reduce throughput on Apple, despite reported 2–6× gains on server GPUs. CUDA kernel designs transfer only partially: public Metal lacks asynchronous copy, and batch-one decode makes several serving-oriented mechanisms inapplicable, although transferable components yield up to 1.20× decode throughput. Quantization gains vary by machine. Across four families, a reported speedup is not a portable property of a technique alone but a joint function of technique, runtime, hardware contract, and workload regime.
1.75×
KV-cache pruning can reach up to 1.75× end-to-end throughput on Apple but can reduce throughput on desktop NVIDIA GPUs.
2–6×
All evaluated speculative decoding methods reduce throughput on Apple, despite reported 2–6× gains on server GPUs.
1.20×
Transferable components of CUDA kernel designs yield up to 1.20× decode throughput.
Introduction
Local LLM inference increasingly happens on Apple machines: llama.cpp treats M-series chips as a first-class target, Apple ships MLX for unified-memory inference, and recent studies describe M-series unified memory as the practical consumer path to large open-weight models. The optimizations that would make this inference fast, however, carry no evidence for it: none of the methods we study - KV-cache management, speculative decoding, the FlashAttention/FlashInfer, Marlin, and MegaBlocks kernels, and post-training quantization - reports a measurement on Apple hardware in its original publication.
An optimization that reduces memory traffic or computation may still fail to improve end-to-end throughput if it targets a secondary cost, adds runtime overhead, or relies on hardware features or workload conditions that are absent from local inference. This motivates our central question: Which LLM inference optimizations remain effective on Apple Silicon, and why?
Our experiments focus on the single-request, unified-memory regime that exists at consumer scale, and we evaluate the bundle users actually deploy: hardware, runtime, and workload together. We cover four Apple devices across three GPU generations, dense and hybrid architectures, and matched RTX 4090 controls where configurations can be aligned.
| Optimization | Methods | Devices | Key finding |
|---|---|---|---|
| KV Cache | Prefix reuse; KVPress pruning (ExpectedAttention, FastKVzip, and 4 others) | M2 Pro, M3 Ultra, M4 Pro; RTX 4090 | Prefix reuse: −57% wall time at turn 9. Pruning: speedup of up to 1.75× on Apple Silicon vs up to 1.02× on RTX 4090. |
| Speculative Decoding | Classic, N-gram, Cascade-n-gram, EAGLE-3, DFlash, DDTree | M2 Pro, M3 Ultra | No method achieved a net speedup. Reported 2–6× gains on NVIDIA GPUs do not generally transfer to different Apple GPUs. |
| Metal Kernels | Marlin (W4A16 GEMM), FlashInfer (split-KV attention), MoE Gather | M2, M2 Pro, M3 Ultra, M4 Pro; RTX 4090 | MoE Gather: +3–12% universally. FlashInfer: +20% on M4 Pro, −3% on M2. Marlin: effective at batch >1, negligible at M=1. |
| Quantization | GGUF/llama.cpp; MLX (affine, DWQ); RTN, QuaRot+RTN, QuaRot+GPTQ | M2 Pro, M3 Ultra; RTX 4090 | MLX native affine/DWQ is the best combined trade-off. QuaRot does not pay off on Apple Silicon without low-precision activation kernels. |
KV Cache
KV caching reduces autoregressive decoding cost by retaining the keys and values computed for preceding tokens. This avoids repeated computation but introduces memory-footprint and bandwidth costs that increase with context length.
Cache growth and persistent prefix reuse
The persistent cache footprint depends on attention architecture and retention policy rather than parameter count alone. At batch size one, it can be expressed as
Cache growth - architecture-derived persistent cache state (bfloat16, MLX-LM layout)
- Qwen3-4B
- Gemma-3-4B, LFM2.5-2.6B, Granite-H-Micro
View data
| Context length (× 1,024 tokens) | Qwen3-4B | Gemma-3-4B | LFM2.5-2.6B | Granite-H-Micro |
|---|---|---|---|---|
| 1 | 144 MiB | 136 MiB | 16 MiB | 45 MiB |
| 4 | 576 MiB | 196 MiB | 64 MiB | 69 MiB |
| 8 | 1,152 MiB | 276 MiB | 128 MiB | 101 MiB |
| 16 | 2,304 MiB | 436 MiB | 256 MiB | 165 MiB |
| 24 | 3,456 MiB | 596 MiB | 384 MiB | 229 MiB |
| 32 | 4,608 MiB | 756 MiB | 512 MiB | 293 MiB |
Qwen3-4B contains 36 global-attention blocks with eight 128-dimensional KV heads, giving 144 MiB at 1,024 tokens and 4,608 MiB at 32,768 tokens. The hybrid models retain less token-indexed attention state. Fewer attention blocks reduce cache growth and the state eligible for pruning.
Persistent caching extends reuse beyond a single response by retaining the longest unchanged prefix across requests, avoiding repeated prefill. In the MLX-LM conversation, wall time at turns 3, 5, and 9 falls by 33%, 42%, and 57%, respectively. These savings reflect avoided prefill rather than increased steady-state decoding throughput.
Dense-model pruning across devices
At 128–256 generated tokens, SnapKV and ExpectedAttention achieve 1.26–1.75× speedup on Apple Silicon, versus 0.45–1.01× on the RTX 4090; FastKVzip remains below baseline on all devices (0.49–0.74×). At a 4,096-token context and 256-token output, ExpectedAttention attains 1.75×, 1.61×, and 1.72× on M2 Pro, M3 Ultra, and M4 Pro, respectively, but 0.98× on the RTX 4090. At a 512-token context, the maximum Apple speedup is only 1.24×.
Matched end-to-end speedup for Qwen3-4B-Instruct-2507 in bfloat16, 256 generated tokens
| 512-token context | 4,096-token context | |||||||
|---|---|---|---|---|---|---|---|---|
| Method (compression) | M2 Pro | M3 Ultra | M4 Pro | RTX 4090 | M2 Pro | M3 Ultra | M4 Pro | RTX 4090 |
| SnapKV (0.5) | 1.13 | 1.11 | 1.16 | 0.96 | 1.42 | 1.36 | 1.41 | 1.01 |
| KNorm (0.5) | 1.11 | 1.11 | 1.13 | 0.94 | 1.50 | 1.37 | 1.38 | 1.00 |
| ExpectedAttention (0.75) | 1.23 | 1.16 | 1.24 | 0.97 | 1.75 | 1.61 | 1.72 | 0.98 |
| StreamingLLM (0.5) | 1.12 | 1.11 | 1.13 | 0.98 | 1.47 | 1.37 | 1.38 | 1.02 |
| TOVA (0.5) | 1.11 | 1.10 | 1.13 | 0.94 | 1.48 | 1.37 | 1.41 | 1.01 |
| FastKVzip (0.75) | 0.56 | 0.28 | 0.52 | 0.67 | 0.68 | 0.49 | 0.71 | 0.58 |
- Slower than baseline (below 1.00)
- Faster, up to 1.75
View data
| Method (compression) | M2 Pro (512-token context) | M3 Ultra (512-token context) | M4 Pro (512-token context) | RTX 4090 (512-token context) | M2 Pro (4,096-token context) | M3 Ultra (4,096-token context) | M4 Pro (4,096-token context) | RTX 4090 (4,096-token context) |
|---|---|---|---|---|---|---|---|---|
| SnapKV (0.5) | 1.13 | 1.11 | 1.16 | 0.96 | 1.42 | 1.36 | 1.41 | 1.01 |
| KNorm (0.5) | 1.11 | 1.11 | 1.13 | 0.94 | 1.50 | 1.37 | 1.38 | 1.00 |
| ExpectedAttention (0.75) | 1.23 | 1.16 | 1.24 | 0.97 | 1.75 | 1.61 | 1.72 | 0.98 |
| StreamingLLM (0.5) | 1.12 | 1.11 | 1.13 | 0.98 | 1.47 | 1.37 | 1.38 | 1.02 |
| TOVA (0.5) | 1.11 | 1.10 | 1.13 | 0.94 | 1.48 | 1.37 | 1.41 | 1.01 |
| FastKVzip (0.75) | 0.56 | 0.28 | 0.52 | 0.67 | 0.68 | 0.49 | 0.71 | 0.58 |
Persistent prefix reuse yields consistent latency reductions by eliminating redundant prefill without token-selection overhead. Pruning gains depend on context and output length, device, and the fraction of inference cost attributable to attention KV state. In the matched Qwen workloads, the largest gains occur on Apple Silicon, while the RTX 4090 remains close to baseline or slows down; FastKVzip slows execution on both platforms. Neither compression ratio nor processor generation alone predicts end-to-end benefit, and reduced KV state does not ensure lower process-level peak memory.
Speculative Decoding
Speculative decoding accelerates autoregressive generation by separating token proposal from verification. A cheaper drafting procedure proposes a sequence of candidate tokens, which the target model evaluates in a single forward pass.
Matched Apple Silicon results
We evaluate Classic, N-gram, Cascade-n-gram, EAGLE-3, DFlash, and DDTree on M2 Pro and M3 Ultra devices with a 4-bit Qwen3-4B target, representative of local inference. All six methods use the same 100 conversational prompts, batch size and concurrency of one, and greedy generation. Prompts average 1,776 tokens, whereas baseline completions average approximately 54–62 tokens.
No configuration in this matched comparison improves E2E throughput. With the 4-bit target, all six methods also reduce decode throughput: losses range from 25.2% to 69.8% on M2 Pro and from 8.9% to 61.7% on M3 Ultra. In the 4-bit M3 Ultra setting, DFlash accepts a mean of 2.57 draft tokens per round and DDTree 3.45. Nevertheless, DFlash reduces decode throughput from 156.7 to 60.0 tokens/s, and DDTree reaches only 61.4 tokens/s. An approximate decode speedup is
Decoding accelerates only when the total round cost is smaller than
Factors limiting speedup
- Workload predictability. The published block-drafting results show substantially less amortizable work on chat than on mathematical reasoning. Our conversational workload is consistent with this less predictable regime.
- Small and quantized targets. A 4-bit 4B target usual for the local on-device inference already has inexpensive autoregressive steps, reducing the work that a speculative round can replace.
- Backend-dependent verification. Block or tree preparation, custom attention masks, verifier traversal, KV-cache updates, and synchronization can add costs that do not scale with single-token target throughput.
- Non-accelerated prefill. The long prompts and short completions make prefill a substantial part of E2E latency. If baseline prefill occupies a fraction
of total latency, unchanged prefill imposes the idealized ceiling , even with cost-free decode.
Configuration-specific positive cases
Supplementary experiments identify configurations with positive decode or E2E changes:
| Method | Device | Target | E2E | Decode |
|---|---|---|---|---|
| Classic | M2 Pro | Qwen3-4B-BF16 | −13.8% | +1.1% |
| Classic | M4 Pro | Qwen3-4B-4bit | −5.0% | +10.0% |
| N-gram | M3 Max | Qwen3-4B-4bit | +2.0% | +13.0% |
| Cascade-n-gram | M3 Max | Qwen3-4B-4bit | −11.6% | +7.6% |
| DFlash | M3 Pro | Qwen3-4B-BF16 | +9.2% | +25.0% |
| DDTree | M3 Pro | Qwen3-4B-BF16 | +15.5% | +40.4% |
Speculative decoding can accelerate particular Apple Silicon configurations, but the published server-GPU gains do not transfer robustly to the principal local chat benchmark. Its benefit depends jointly on accepted work, target-step cost, the complete draft–verify path, and prefill's share of latency.
Metal Kernels
CUDA kernels exploit asynchronous copies, matrix units and cooperative grids; MLX on Apple Silicon exposes a different contract. We test which mechanisms from Marlin, FlashInfer and grouped MoE GEMM survive in single-request decode.
The transfer taxonomy
A mechanism that fails to port does so for one of four reasons:
| Label | Meaning |
|---|---|
| TRANSFERS | The idea and its mechanism both exist on Metal. Port it. |
| HW-GAP | The idea is sound but the hardware or language offers no equivalent instruction. Nothing you write closes this. Example: cp.async. |
| RUNTIME | The overhead the technique removes has already been removed by the runtime, so the ceiling is near zero before you start. Example: kernel-launch amortisation under MLX's lazy graph. |
| N/A | The technique's precondition does not hold in our regime. Not a platform limitation at all - the same code on the same hardware would need it if the regime changed. Example: sorting tokens by expert when the batch is one token. |
Marlin: weight-only 4-bit GEMM
Register-only reconstruction, group factoring and repacking transfer; unpacking needs different arithmetic. The matrix unit is N/A at M = 1, and public MSL makes asynchronous prefetch a HW-GAP. At M = 1 the kernel measures 0.95×–1.00× against MLX qmv and 0.993× end to end. The picture changes at M = 2–8: dequantising each block once into registers and reusing it across activation rows yields 1.09×–1.20× on all tested Qwen3-4B and Qwen3-8B transformer-layer shapes.
FlashInfer: decode attention
Online softmax, split-KV and zero-copy cache views transfer. Paging and load-balanced planning require concurrent serving and are N/A; asynchronous prefetch is a HW-GAP. End to end, the gain reaches 1.20× on M4 Pro but falls to 0.97× on M2.
FlashInfer decode-throughput speed-up across context lengths (Qwen3-4B)
- M4 Pro
- M2
View data
| Context length (tokens) | M4 Pro | M2 |
|---|---|---|
| 1k | 1.02× | 1.02× |
| 5k | 1.08× | 0.99× |
| 10k | 1.15× | 0.98× |
| 16k | 1.20× | 0.97× |
MoE gather kernel
All four mechanisms assume many rows per expert. At batch 1 there is nothing to sort, no group larger than one row, no tile quantisation to avoid and no load-balancing problem. Against mx.gather_qmm, the kernel reaches 1.05× on gate/up and 1.23× on down; patching 66 SwitchLinear layers yields 1.03×–1.12× end to end across machines. The gain comes from removing machinery batch-1 decode does not need, not from recreating grouped GEMM.
Exploratory megakernel case study
The prototype fused eleven operations into one bit-faithful dispatch; with the barrier disabled to isolate compute, it measured 174 µs against 96 µs for the stock chain, a 1.8× slowdown. In this Apple/MLX configuration, dispatch boundaries are nearly free and byte traffic dominates.
Conclusions
Portability depends on whether the bottleneck and workload survive the move. Marlin's weight reuse helps at M > 1, making it useful for batching and speculative verification, but not for bandwidth-bound M = 1. FlashInfer's split-KV is more relevant on newer Apple GPUs that can repay its extra dispatch and partials traffic. MoE gather is positive on every tested machine and context because it removes machinery unused at batch 1. Public Metal's missing asynchronous copy remains a common gap, and the RTX controls show that CUDA outcomes do not reliably predict results on Apple Silicon.
Quantization
Post-training quantization reduces the model's memory footprint by lowering weight precision after training, trading a small loss in quality for faster inference and lower hardware requirements. We evaluate weight-only post-training quantization of LFM2.5-8B-A1B across three independent toolchains: GGUF/llama.cpp, MLX, and a custom llm-compressor pipeline covering RTN, QuaRot+RTN, and QuaRot+GPTQ, all served on MLX kernels.
Latency and throughput
Decode gains depend on bandwidth headroom. On both chips, decode throughput rises as bit-width drops, but the extent depends on how memory-bound the chip already is. M2 Pro cannot hold the BF16 checkpoint at all, so here quantization is a precondition for running the model rather than an optimization.
Prefill gains less from quantization than decode. Moving from 8- to 4-bit raises MLX-native prefill by only 2–7% on M2 Pro and 4–6% on M3 Ultra, against 20–64% for decode, consistent with prefill being compute-bound rather than bandwidth-bound.
The best runtime depends on the chip. On M2 Pro, GGUF-Q4_0 is the clear choice. On M3 Ultra, the two runtimes are close: GGUF-Q4_0 leads the MLX-native formats on TTFT and prefill, while MLX-W4-Affine leads on decode. In our weight-only setting, QuaRot's rotations are pure overhead: it costs 2.4–4.2% of decode and about 1% of long-scenario prefill relative to RTN-Asym.
Cross-platform comparison
Relative decode gains differ across devices. Normalized to each device's own Q8_0, Q4_0 improves decode by +63.7% on M2 Pro, +38.2% on the RTX 4090, and +20.4% on M3 Ultra. M2 Pro, the most bandwidth-constrained device, benefits the most.
The k-quant prefill slowdown is specific to Apple Silicon. On the RTX 4090, Q5_K_M prefill is essentially flat relative to Q8_0 (+0.9%), whereas it falls by 17.3% on M2 Pro and 9.8% on M3 Ultra.
Conclusions
As with the kernel results, a reported quantization benefit does not automatically carry over across platforms. The same Q8_0-to-Q4_0 step speeds up GGUF decode by 1.64× on M2 Pro but only 1.20× on M3 Ultra, and the k-quant prefill slowdown on both Apple chips is largely absent on the RTX 4090. Without low-precision activation kernels, QuaRot's rotations are pure overhead.
Discussion and Conclusion
Across four optimization families, no single verdict holds uniformly. Each technique's fate on Apple Silicon depends on a specific, identifiable mismatch between its assumptions and our batch-1, unified-memory regime.
KV-cache pruning helps only when the selection cost is low relative to the attention it removes. Profit grows with output length, context, and architecture, and reverses on NVIDIA GPUs. Speculative decoding produced no standardized E2E gain in the principal long-prompt chat benchmark. Draft–verify overhead and prefill penalties offset its decode-side benefits, although separate configurations showed positive results. Metal kernel ports succeed only where the workload matches their design. Marlin's weight reuse pays off only once more than one token decodes at a time, useful for batching and draft verification but not single-token chat; FlashInfer's split-KV attention earns back its overhead only on newer Apple GPUs; and the MoE gather helps everywhere, since it removes scheduling work batch-one decoding never needed. Quantization gains reflect how memory-bound the chip already is, rather than bit-width itself. Neither GGUF nor MLX is better in all scenarios. Generally, GGUF is faster on prefill, but MLX has better decode speed, so the best choice depends on which phase matters more for the workload.
The unifying result is that a reported speedup is not a portable property of a technique alone, but a joint function of technique, runtime, hardware contract, and workload regime. Practitioners porting CUDA-validated techniques to Apple Silicon should identify which bottleneck a mechanism actually removes before expecting a number.
Limitations. Our KV-pruning lacks task-quality evaluation, and the megakernel study covers a single model on a single chip. Extending this analysis to multi-request serving, where batching gives several currently idle mechanisms something to act on, is left to future work.
Citation
@misc{yemets-etal-2026-neurips-odi-apple-silicon,
author = {Kyrylo Yemets and Andrii Yaroshevych and Victor Muryn and Ostap Khomenko and Pavlo Kryven and Maksym Kmet and Maksym Shamrai and Mariya Hirna},
title = {Evaluating the Applicability of {LLM} Inference Optimizations on {Apple Silicon}},
year = {2026},
eprint = {2610.31479},
archivePrefix = {arXiv},
primaryClass = {cs.LG},
url = {https://arxiv.org/abs/2610.31479},
note = {\emph{Accepted to the NeurIPS 2026 Workshop on On-Device Intelligence (ODI).}},
}