**From CUDA to MLX: How K-Search Brings Decades of Kernel Expertise to Apple Silicon**
Modern AI workloads are driving rapid innovation in hardware, with an expanding ecosystem of chips from NVIDIA, Apple, and specialized accelerators. Alongside this shift, AI coding tools have drastically shortened the time needed to prototype kernels. Yet, high-performing GPU kernels remain a bottleneck—writing efficient ones requires years of expertise, and transferring these kernels between architectures often means reinventing the wheel. IBM Research bridges this gap by extending K-Search, an evolutionary kernel search framework from Berkeley Sky Lab, with a CUDA-to-MLX translation layer that brings decades of NVIDIA kernel expertise to Apple Silicon. This approach achieves near-native performance and dramatic speedups, without starting from scratch.
—
### Why MLX?
Apple’s MLX framework has seen rapid adoption, running AI models locally on Macs with Apple Silicon. Its unified memory model is particularly attractive for mid-sized models, reducing reliance on cloud inference. However, MLX lacks many performance-critical optimizations found in mature CUDA ecosystems—such as paged attention, optimized state-space model (SSM) kernels, and fused multi-head attention. MLX runs models correctly but often leaves performance on the table. This gap motivated the development of a systematic way to transplant CUDA expertise to MLX.
—
### What Is K-Search?
K-Search is an evolutionary framework for optimizing GPU kernels. Given a naive kernel and hardware specifications, it iteratively explores optimization strategies guided by an LLM. The LLM acts as a “GPU kernel performance engineer,” classifying the kernel, analyzing bottlenecks, and proposing candidate optimizations. These candidates are compiled and benchmarked on real hardware; measurements feed back into a persistent search tree, enabling the system to refine promising paths and abandon dead ends.
K-Search relies on a *Spec* document that encodes hardware rules and constraints, preventing invalid code generation. In prior work, K-Search consistently outperformed competing evolutionary methods on CUDA kernels such as FlashAttention. The key question addressed in this work: *Can this CUDA expertise be transferred to non-CUDA platforms like Apple Silicon?*
—
### Building an MLX Backend
Extending K-Search to MLX required more than a direct port. The team built a complete MLX task backend, including:
– Kernel compilation and execution via MLX’s Metal/C++ APIs
– Tailored kernel generator prompts for Metal
– Integration with MLX’s benchmarking tools
This backend enables K-Search to operate natively on Apple Silicon, forming the foundation for cross-platform knowledge transfer.
—
### Translating CUDA Expertise to MLX
The central challenge was transferring decades of optimization knowledge rather than copying kernels instruction-by-instruction. The translation layer includes:
– **Concept mapping tables:** Systematic mappings from CUDA primitives to MLX/Metal equivalents, with hardware constraints. For example:
– `__shared__` → Metal threadgroup memory (32 KB limit vs. 48 KB on NVIDIA)
– `warp_reduce` → Matrix Multiply-Accumulate (MMA) operations
– `__syncthreads()` → `threadgroup_barrier`
– Bandwidth assumptions calibrated to Apple Silicon (e.g., M3 Max vs. H100)
– **MLX-specific patterns:** Optimizations with no direct CUDA analog, such as register-based reductions using `simd_shuffle_xor` and the “exp2 trick” for faster softmax.
– **Reusable assertions:** High-level properties that must be preserved during optimization, guiding the evolutionary search without hardcoding implementation.
This structured approach ensures generated kernels are both valid and performant.
—
### Matching Expert Kernel Performance: Attention
The attention kernel evaluation compared three configurations:
1. A naive MLX baseline
2. K-Search with no additional context
3. K-Search with the full CUDA translation layer
Results showed a dramatic improvement from the translation layer, raising performance from 0.26× to 0.97× of Apple’s native MLX attention kernel. With full context, the evolved kernel independently discovered key FlashAttention-2 optimizations:
– Threadgroup memory tiling
– Online softmax
– K-transposition for memory access
– The exp2 trick, replacing $e^x$ with $exp_2(x cdot log_2(e))$ to leverage Apple’s fast `exp2()` hardware instruction
—
### A 20× Faster Prefill: The Mamba SSM Kernel
Beyond attention, K-Search was applied to Mamba’s state-space model kernel. On an M1 Max, the evolved MLX kernel achieved up to **20× prefill speedup** over the community MLX implementation. This gain stems from a parallel prefix scan, which exploits the associative property of state transitions:
[
(a_2, b_2) circ (a_1, b_1) = (a_2 a_1, a_2 b_1 + b_2)
]
While the community kernel processed tokens sequentially, the evolved kernel applied a parallel scan, reducing complexity from $O(N)$ to $O(log N)$. Decode performance remained comparable, as expected, since decode operates token-by-token.
—
### What’s Next?
The work demonstrates that AI-driven kernel search, grounded in structured cross-platform translation, can deliver near-expert performance on Apple Silicon. Future directions include:
– Expanding to new architectures, including IBM Spyre AIU
– Adding support for more kernels (e.g., paged attention, fused MoE)
– Automating translation context within the K-Search loop
—
### FAQs
**Q: What is K-Search?**
A: K-Search is an evolutionary framework that uses an LLM to explore kernel optimizations iteratively. It maintains a tree of hypotheses, evaluates candidates on real hardware, and refines promising paths based on measured performance.
**Q: How does the CUDA-to-MLX translation layer work?**
A: It maps CUDA concepts—such as shared memory, warp-level primitives, and synchronization—to their MLX/Metal equivalents with hardware constraints. It also provides MLX-specific optimization patterns and correctness assertions to guide the search.
**Q: Which kernels were optimized?**
A: The attention kernel and the Mamba SSM kernel were the primary targets, though the approach is not architecture-specific.
**Q: How much performance improvement was observed?**
A: Attention reached 97% of native MLX kernel performance, and Mamba prefill achieved up to 20× speedup over community implementations.
**Q: Is this approach limited to Apple Silicon or MLX?**
A: No. While the backend targets MLX, the translation framework is general and can apply to any non-CUDA platform where CUDA expertise needs to be transferred.
—
### Conclusion
By combining K-Search with a structured CUDA-to-MLX translation layer, IBM Research successfully transferred decades of GPU kernel expertise to Apple Silicon. The results—near-expert attention kernels and a 20× prefill speedup for SSMs—highlight the power of evolutionary search guided by cross-platform optimization knowledge. This approach offers a scalable path toward high-performance AI kernels across diverse hardware, reducing reliance from-scratch tuning and accelerating deployment on emerging architectures.
