PaperlayerView exact revision on arXiv

Source-linked AI summary

Evaluating the Applicability of LLM Inference Optimizations on Apple Silicon

Kyrylo Yemets, Andrii Yaroshevych, Victor Muryn, Ostap Khomenko, Pavlo Kryven, Maksym Kmet, Maksym Shamrai, Mariya Hirna

arXiv:2610.31479v1cs.LGcs.DC

TL;DR

This paper asks which inference optimizations designed largely for server GPUs remain effective in local Apple Silicon inference. It evaluates four optimization families on Apple machines and matched RTX 4090 controls where possible. Results vary by technique and setting, with no uniform verdict across the families.

  • Problem

    The studied inference optimizations lack measurements on Apple hardware, leaving their effectiveness in local Apple Silicon inference uncertain.

    Source: §1 · p. 1

  • Method

    The paper evaluates four optimization families in single-request unified-memory inference across Apple devices, using matched RTX 4090 controls where configurations align.

    Source: §1 · p. 2

  • Results

    No single verdict holds across the four optimization families, and no configuration in the matched speculative-decoding comparison improves end-to-end throughput.

    Source: §2.2.1 · p. 7, §3 · p. 27

  • Takeaways & Limitations

    Portability depends on whether an optimization’s bottleneck and workload survive the move; CUDA outcomes do not reliably predict Apple results.

    Source: §2.3.7 · p. 21

    The megakernel interpretation is scoped to public Metal/MLX designs requiring grid-wide barriers, not an impossibility result for other designs or APIs.

    Source: P1: what a dispatch boundary actually… · p. 18

Abstract

from arXiv · show

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 times 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 times 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 times 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 Introduction

Although local LLM inference is increasingly common on Apple machines, the studied optimizations had not been measured on Apple hardware in their original publications. This work isolates optimization effects across Apple devices and matched NVIDIA controls to assess when those effects transfer.

1 Introduction

  • The studied optimization publications report no measurements on Apple hardware, despite growing local inference use on Apple machines.The gap spans KV-cache management, speculative decoding, GPU kernels, and post-training quantization.Source: §1 · p. 1
  • The study asks which LLM inference optimizations remain effective on Apple Silicon and why.Reducing memory traffic or computation may not improve end-to-end throughput when an optimization targets a secondary cost, adds runtime overhead, or assumes absent hardware features or workload conditions.Source: §1 · p. 1
  • Experiments evaluate four Apple devices across three GPU generations, with matched RTX 4090 controls where configurations align.The focus is single-request unified-memory inference, including dense and hybrid architectures.Source: §1 · p. 1
  • Unlike prior Apple-Silicon studies of models or complete software stacks, this work isolates individual optimizations and measures kernel-level and end-to-end effects.It ports representative CUDA mechanisms to Metal/MLX and analyzes when gains persist, disappear, or reverse across hardware, runtime, and workload differences.Source: §1 · p. 2

2 Methods and Results

This section evaluates KV-cache pruning, speculative decoding, kernel porting, and quantization across Apple hardware, runtimes, and workload conditions. Across these methods, optimization gains depend on the platform and workload, so CUDA results do not reliably predict Apple performance.

2.1.2 Cache growth and persistent prefix reuse

  • Qwen3-4B cache grows from 144 MiB at 1,024 tokens to 4,608 MiB at 32,768, while persistent prefix reuse cuts wall time by 33%, 42%, and 57% at turns 3, 5, and 9.Cache growth depends on attention architecture and retention policy, not parameter count alone; prefix-reuse savings come from avoided prefill, not higher steady-state decode throughput.Source: §2.1.2 · p. 3, Figure 1 · p. 3

2.1.3 Dense-model pruning across devices

  • At a 4,096-token context and 256-token output, Qwen pruning reaches 1.75× on M2 Pro, versus 0.98× on RTX 4090; FastKVzip stays below baseline across the grid.The matched batch-one experiments compare Qwen3-4B-Instruct-2507 and LFM2.5-2.6B across Apple devices, with RTX 4090 included for Qwen; tested methods use compression ratios of 0.5 or 0.75.Source: §2.1.1 · pp. 2–3, §2.1.3 · p. 4, Figure 3 · p. 5

2.1.4 Hybrid-model pruning and peak memory

  • At 4,096 tokens and 256 output tokens, LFM2.5-2.6B pruning reaches at most 1.25×, while M4 Pro peak-memory savings range from 20.4% to 25.6% for Qwen.At 512-token contexts, pruning increases peak memory by up to 155.5% for Qwen and 48.5% for LFM, consistent with temporary scoring and compaction allocations outweighing cache savings.Source: §2.1.4 · p. 4, Figure 4 · p. 5, Figure 5 · p. 6

2.2.1 Matched Apple Silicon results

  • With a 4-bit Qwen3-4B target, all six speculative methods reduce decode throughput and none improves end-to-end throughput on either M2 Pro or M3 Ultra.Decode losses range from 25.2% to 69.8% on M2 Pro and from 8.9% to 61.7% on M3 Ultra; BF16 also yields no end-to-end gain, with M2 Pro Classic the sole decode-positive case at +1.1%.Source: §2.2.1 · p. 7, Figure 6 · p. 7

2.2.2 Factors limiting speedup

  • Speculation pays only when accepted work offsets the full draft–verify cost, and short, predictable-target workloads and non-accelerated prefill constrain gains.The tested 4-bit 4B target has inexpensive autoregressive steps, while drafting and verification overhead can persist; long prompts and short completions also limit end-to-end benefit, and other length regimes were not directly evaluated.Source: §2.2.1 · p. 8, §2.2.2 · pp. 8–9

2.2.3 Configuration-specific positive cases

  • DFlash and DDTree on M3 Pro show the largest observed end-to-end gains, 9.2% and 15.5%, but these configuration-specific positives do not establish a controlled method or processor ranking.The results support a cost trade-off involving accepted work, target-step cost, draft–verify path, and prefill share; gains do not persist across other device tests.Source: §2.2.3 · p. 9, Table 2 · p. 9

2.3.1 The transfer taxonomy

  • The kernel-transfer taxonomy distinguishes TRANSFERS, HW-GAP, RUNTIME, and N/A: Metal may lack an equivalent instruction, the runtime may remove an overhead, or batch-one may violate a technique’s precondition.The evaluation tests CUDA mechanisms from Marlin, FlashInfer, and grouped MoE GEMM in single-request decode, with matched RTX 4090 controls where possible.Source: §2.3.1 · pp. 9–10, §2.3 · p. 9
  • The transfer taxonomy treats RUNTIME as a measure-before-kernel instruction and N/A as workload-dependent deferral, not refutation.For ported techniques, the authors argue from the distribution of transfer verdicts rather than the speed-up produced by a port.Source: §2.3.1 · p. 10

2.3.2 Marlin: weight-only 4-bit GEMM

  • Marlin avoids erasing W4A16's bandwidth savings by reconstructing weights in registers and consuming them without writing expanded values to memory.The Metal port reaches the unpacking idea through exact power-of-two prescaling rather than NVIDIA's lop3 instruction; grouping enables the affine scale and bias to be applied once per group.Source: §2.3.2 · p. 10, Mechanisms · p. 10

Mechanisms

  • At M = 2–8, the Marlin-derived kernel delivers 1.09×–1.20× on tested Qwen3-4B and Qwen3-8B transformer-layer shapes.At M = 1 it measures 0.95×–1.00× against MLX qmv and 0.993× end to end; register-only reconstruction, group factoring, and repacking transfer, while asynchronous prefetch is inaccessible through public MSL and the matrix unit is N/A at M = 1.Source: Mechanisms · pp. 10–12, Table 3 · p. 11

2.3.3 FlashInfer: decode attention

  • FlashInfer's online softmax supports one-pass decode attention, while split-KV partitions the cache to restore batch-one parallelism.Paged caches and host-side chunk load balancing are serving-oriented mechanisms: MLX uses a growing contiguous cache and batch one has nothing to load-balance. The Metal kernel reads cache strides without copying the view.Source: §2.3.3 · pp. 12–13, Table 5 · p. 13
  • FlashInfer reaches 1.20× end-to-end decode throughput on M4 Pro but 0.97× on M2.The same shader measured 0.63× at N = 256 and 0.73× at N = 1024 through a Python wrapper, versus 1.06× at both lengths as a native MLX primitive, showing the wrapper overhead caused the short-context loss. Split-KV can add fixed dispatch and partial-result costs when a query already occupies the GPU, so it is enabled per architecture from measured results.Source: Figure 7 · p. 14, §2.3.3 · p. 14, Table 6 · p. 14

Follow-up experiments

  • The MoE gather transfer table rates sorting, grouped scheduling, block sparsity, and block fusion N/A for batch-one decode, where their many-row preconditions do not hold.A follow-up down-projection-plus-routing-reduction fusion adds 1.2–3.0%; gate/up/SwiGLU fusion measures 0.986×–0.999×, repeated-expert staging 0.76×–0.93×, and prefill grouped GEMM 0.47× against stock on M4 Pro.Source: Mechanisms and their preconditions · p. 15, Follow-up experiments · p. 16

P1: what a dispatch boundary actually costs

  • The megakernel prototype takes 174 µs versus 96 µs for the stock chain, a 1.8× slowdown on one M4 Pro LFM2.5-8B-A1B block.In this MLX configuration, token pipelining hides the 0.54 ms of graph building, while the removable block-boundary cost is 15 µs and forcing phases onto one grid costs 78 µs; public Metal lacks a cooperative-launch and occupancy-query contract. This is a scoped case study, not an impossibility result for megakernels.Source: §2.3.5 · p. 17, P1: what a dispatch boundary actually… · pp. 17–18

2.3.6 RTX 4090 comparison

  • The RTX 4090 controls compare dependent-versus-independent matvec chains and dispatch-boundary cost, not every kernel verdict.Matched runs use PyTorch 2.13 and CUDA 13.2; the experiments control for L2 residency, GQA expansion, and implicit batching to measure single-request decode.Source: §2.3.6 · p. 18
  • The RTX 4090 incurs a 57.4% dependent-chain penalty; CUDA graph capture recovers 91.2% of boundary cost, while 24-stream concurrency does not recover the gap.The matched chain includes interstitial elementwise dispatches; the authors interpret its penalty as a dependency-chain cost, not a pure matvec roofline effect.Source: §2.3.6 · p. 18, Figure 8 · p. 19, Figure 9 · p. 19

Decode attention

  • The RTX 4090 FlashInfer comparison is stack-level rather than hardware-only because its baseline differs from the Apple comparison; at 1k, both kernels remain far below peak bandwidth.The matched speculative-decoding configuration also changes sign across stacks: at k = 7, the RTX 4090 path reaches 0.703× versus 0.99× on Apple and 1.13× with the small-M Metal verifier kernel.Source: Table 13 · p. 20, Table 14 · p. 20

Speculative decoding

  • Two NVIDIA comparisons are withheld: a Marlin layout mismatch invalidated timings, and batch-one MoE gather comparisons were uninformative because both paths were launch-bound.The reported NVIDIA Table 12 comparison therefore uses the stock int4 path; the MoE N/A verdict rests on the Apple-side prefill result.Source: Speculative decoding · p. 20

2.3.7 Conclusions

  • Portability depends on whether the bottleneck and workload survive the move: Marlin weight reuse helps at M > 1, while FlashInfer split-KV is more relevant on newer Apple GPUs.MoE gather is positive across tested machines and contexts because it removes machinery unused at batch 1; public Metal still lacks asynchronous copy.Source: §2.3.7 · p. 21

2.4.1 Latency and throughput

  • On M2 Pro, GGUF-Q4_0 leads overall; on M3 Ultra, GGUF-Q4_0 leads prefill while MLX affine/DWQ variants lead decode.M2 Pro cannot hold the BF16 checkpoint, making quantization a precondition for running the model there. On M3 Ultra, RTN-Asym has the highest E2E throughput.Source: Table 15 · p. 21, Table 16 · p. 22, §2.4.1 · p. 23

2.4.2 Cross-platform comparison

  • Q4_0 decode gains are +63.7% on M2 Pro, +38.2% on RTX 4090, and +20.4% on M3 Ultra; the k-quant prefill slowdown seen on Apple is largely absent on the 4090.The RTX 4090 is fastest at every quantization level; its Q4_0 decode is 469.1 tok/s, versus 232.2 tok/s on M3 Ultra and 129.3 tok/s on M2 Pro.Source: §2.4.2 · pp. 23, 25, Figure 13 · p. 24, Figure 12 · p. 24

2.4.3 Memory footprint

  • Quantized memory falls roughly with bit-width, but RTN/QuaRot peaks at 6,630–6,640 MB versus 5,930–6,302 MB for MLX 4-bit variants, indicating greater transient pipeline overhead.The similar model-memory footprints and comparable gap for plain RTN-Asym suggest the extra overhead comes from checkpoint loading or conversion rather than QuaRot itself.Source: §2.4.3 · p. 25, Table 17 · p. 25

2.4.4 Downstream quality

  • GGUF-Q6_K’s 58.4 average quality matches rather than exceeds BF16 within confidence intervals, while all four QuaRot variants lose 9.7–12.1 IFBench points.The QuaRot IFBench drop exceeds the ±5.6-point confidence interval; without low-precision activation kernels, rotations add overhead and reduce quality in this pipeline.Source: §2.4.4 · p. 26, Table 18 · p. 26, §2.5 · p. 26

3 Discussion and Conclusion

Across the four optimization families, no single verdict holds uniformly: Apple Silicon outcomes depend on how each technique’s assumptions match the batch-1, unified-memory regime. The paper concludes that speedups depend jointly on technique, runtime, hardware contract, and workload regime.

3 Discussion and Conclusion

  • A technique’s speedup is not portable by itself; its outcome depends jointly on technique, runtime, hardware contract, and workload regime.The paper advises checking which bottleneck a mechanism removes before expecting a reported gain to transfer.Source: §3 · p. 27
  • KV-cache pruning helps when selection costs less than the attention it removes, with gains growing by output length, context, and architecture, but reversing on NVIDIA GPUs.The paper also notes that pruning-induced output-quality effects were not evaluated.Source: §3 · p. 27
  • No speculative-decoding method produced a standardized end-to-end gain in the principal long-prompt chat benchmark, as drafting, verification, and prefill costs offset decode benefits.Separate configurations showed positive results.Source: §3 · p. 27
  • Metal kernel ports help only when workload conditions fit their design: Marlin reuse needs multiple decoding tokens, FlashInfer split-KV pays off only on newer Apple GPUs, and MoE gather helps at batch one.The MoE gather removes scheduling work that batch-one decoding does not need.Source: §3 · p. 27
  • Quantization gains reflect how memory-bound a chip already is, while neither GGUF nor MLX wins in every scenario; GGUF is generally faster at prefill and MLX at decode.The better runtime choice therefore depends on which phase matters more for the workload.Source: §3 · p. 27
  • The study does not evaluate KV-pruning task quality, and its megakernel analysis covers one model on one chip; multi-request serving remains future work.The paper identifies batching as a regime that could give currently idle mechanisms something to act on.Source: §3 · p. 27

A Additional Speculative-Decoding Results

The appendix records implementation settings and a supplementary 8B sensitivity experiment. Method review, acceptance and cost analysis, limiting factors, and configuration-specific positive cases are covered in Section 2.2.

A Additional Speculative-Decoding Results

A.1 Implementation and benchmark configuration

The benchmark configuration uses paired autoregressive baselines for speculative methods and holds method-specific settings fixed across two Apple devices and target precisions unless stated otherwise. It specifies the chat workload, method variants, and conventions for interpreting throughput and acceptance.

A.1 Implementation and benchmark configuration

  • The principal comparison uses batch size and concurrency one on M2 Pro and M3 Ultra, with 100 conversational prompts averaging 1,776 prefill tokens.The dataset includes 46 tool-calling and 54 regular-chat samples; generation is greedy with a 512-token limit, one warmup, and three measured repetitions per prompt and method.Source: §A.1 · p. 31
  • Classic uses a 0.6B 4-bit Qwen3 drafter with a four-token draft width, the fastest setting among those tested.The same drafter is used for both target precisions.Source: §A.1 · p. 31
  • N-gram proposes two tokens from the longest matching stored suffix, while Cascade-n-gram uses the same table to speculate for the 0.6B drafter before target-level drafting.The cascade keeps a two-token target-level draft width and evaluates a vertical N-gram cascade, not Chen et al.’s full vertical and horizontal Max-Gram cascade [17].Source: §A.1 · p. 31
  • EAGLE-3 uses a four-bit converted community head, drafts one token per step, and takes hidden states from target layers 2, 18, and 33.This local configuration differs from the wider dynamic-tree settings used for the published speedups.Source: §A.1 · p. 31
  • DFlash and DDTree share a BF16 block drafter; DFlash uses a 16-position block including the target-produced root, while DDTree budgets 24 non-root tree nodes.Both use greedy verification.Source: §A.1 · p. 31
  • Standardized end-to-end throughput uses each method’s paired autoregressive baseline token workload, so absolute rates across implementations are not a controlled cross-method comparison.AL excludes the target-produced root for DFlash and DDTree, while committed-token count τ includes target-produced progress as applicable; their draft-budget definitions also differ.Source: §A.1 · p. 31

A.2 Additional results for Qwen3-8B-4bit

With the 8B target, relative end-to-end performance improves for most speculative configurations, but none gains throughput. N-gram comes closest on M3 Ultra, and decode effects remain method-dependent.

A.2 Additional results for Qwen3-8B-4bit

  • No speculative method achieves an end-to-end gain with Qwen3-8B-4bit; N-gram comes closest on M3 Ultra at −1.4%.The comparison reports standardized end-to-end, prefill, and decode throughput changes relative to each implementation’s autoregressive baseline on M2 Pro and M3 Ultra.Source: Table 19 · p. 31, §A.2 · p. 32
  • Increasing the target from 4B to 8B improves relative end-to-end performance for most configurations without eliminating drafting, verification, and prefill costs.The larger target provides additional amortization headroom, but no method reaches an end-to-end gain.Source: §A.2 · p. 32
  • Decode changes depend on the method: EAGLE-3 becomes slower relative to baseline on both devices, while DFlash and DDTree modestly reduce their penalties.The workload, decoding protocol, and paired-baseline normalization follow the principal comparison.Source: §A.2 · p. 32

B.0.1 Marlin: 6-bit lm_head

The 6-bit lm_head matvec matched stock MLX, indicating that cheaper unpacking did not improve this rate-capped shape.

B.0.1 Marlin: 6-bit lm_head

  • 835 µs versus 844 µs: the 6-bit tied lm_head matvec measured parity with stock MLX.The tested layout used 4+2 bit planes.Source: §B.0.1 · p. 33
  • MLX sustained 252 GB s−1–264 GB s−1 on this shape at 4, 6 and 8 bits alike.The passage identifies the shape as rate-capped, so cheaper unpacking had nothing to recover.Source: §B.0.1 · p. 33
  • The 4+2 bit-plane matvec therefore did not outperform stock MLX on the tested tied lm_head shape.The measured parity is consistent with unpack cost not being the limiting resource.Source: §B.0.1 · p. 33

B.0.2 FlashInfer: split-KV routing calibration

A fixed split-KV threshold misrouted some Apple workloads and cost performance on M2, so the port uses measured per-architecture routing.

B.0.2 FlashInfer: split-KV routing calibration

  • 1–3%: the reference N ≥1024 threshold cost performance on M2 by forcing two-pass execution where stock MLX stayed single-pass.At 5–7k context, that threshold also under-split by 2×.Source: §B.0.2 · p. 33
  • The router uses a measured per-architecture table rather than applying the same split-KV threshold to every GPU.For an unmeasured machine, it inherits stock behavior instead of a heuristic tuned on different hardware.Source: §B.0.2 · p. 33
  • The reference heuristic enables split-KV at N ≥1024, while the MLX routing path uses gqa > 1 and N ≥4096 outside the listed device classes.The routing decision is resolved once at import.Source: §B.0.2 · p. 33

B.0.3 Megakernel: cooperation, occupancy and geometry

The exploratory megakernel probes on one M4 Pro block found that Metal can implement a grid barrier with atomics, but residency and cross-threadgroup visibility constrain its safe use.

B.0.3 Megakernel: cooperation, occupancy and geometry

  • The probes complete the exploratory case study of §2.3.5, covering one LFM2.5-8B-A1B MoE block on one M4 Pro under MLX.Source: §B.0.3 · p. 33
  • 96% of ordinary device stores were stale across threadgroups, while relaxed atomics had 0 stale reads in 41.3 × 10^6 reads.The tested cross-phase handoffs therefore had to round-trip through relaxed atomics.Source: §B.0.3 · p. 33
  • Public Metal has no cooperative launch; a hand-rolled grid barrier worked, but requires every threadgroup to be resident or spinning groups can deadlock.The listing notes there is no API to determine residency in advance.Source: §B.0.3 · p. 33

P3: the same kernel, one register-footprint apart

The megakernel probes found that fixed-grid fusion can lose more to poor grid geometry than it recovers by removing dispatch boundaries, while register pressure made convergence unpredictable.

P3: the same kernel, one register-footprint apart

  • Changing only per-thread register footprint separated a convergent variant from one that timed out in a silent 30 ms spin.The compiler does not report the footprint, and no occupancy query predicts the outcome.Source: P3: the same kernel, one… · p. 34
  • 174 µs versus 96 µs: the one-dispatch megakernel compute path was slower than the stock 11-op chain.Effective rate was 124 GB s−1 for the megakernel versus 225 GB s−1 over 22 MB for stock.Source: P3: the same kernel, one… · p. 34
  • 78 µs: shape-optimal grids were worth five times the 15 µs recovered by removing every boundary on this block.The fused phase-A grid streamed 14.7 MB of gate and up weights at one-third the rate of the same code in its own dispatch.Source: P3: the same kernel, one… · p. 34
  • A Mirage-style interpreter would retain combined worst-case register allocation and cannot prefetch next-instruction weights without a public cp.async-class mechanism.The passage presents this as a constraint for the evaluated workload.Source: P3: the same kernel, one… · p. 34

C Full Quantization Benchmark Results

This section reports supplemental quantization benchmark measurements not shown in the main text.

C Full Quantization Benchmark Results

C.1 Latency and throughput

Table 20 reports latency and throughput for LFM2.5-8B-A1B on the M2 Pro in the short scenario.

C.1 Latency and throughput

C.2 Cross-platform comparison

The supplied results describe RTX 4090 GGUF quantization and M3 Ultra quantization benchmarks, without stating their comparative outcomes.

C.2 Cross-platform comparison

  • Figure 14 plots RTX 4090 peak memory footprint versus throughput for GGUF quantization of LFM2.5-8B-A1B.Source: §C.2 · p. 35, Figure 14 · p. 36
  • Table 21 reports short-scenario latency and throughput for LFM2.5-8B-A1B on M3 Ultra.The scenario uses a 512-token prompt and 256-token generation.Source: Table 21 · p. 36
  • Table 21 covers GGUF, MLX, and llm-compressor (RTN/QuaRot), with a separate BF16 baseline for each runtime.Source: Table 21 · p. 36

Related papers

Loading 2610.31479v1…