K-Search: Transferring CUDA Kernel Expertise to Apple Silicon MLX

TL;DR

Researchers have extended K-Search, an evolutionary kernel optimization framework, with a structured translation layer that adapts decades of CUDA kernel expertise for Apple Silicon's MLX framework. This approach allows AI to automatically generate high-performance GPU kernels, achieving 0.97x the speed of native MLX Attention kernels and up to a 20x prefill speedup for Mamba SSM kernels compared to community implementations.

The Challenge of Cross-Platform Kernel Optimization

Writing efficient GPU kernels requires years of specialized expertise, and transferring these optimizations from one hardware vendor to another typically requires rediscovering them from scratch. While the CUDA ecosystem possesses extensive hand-tuned implementations for critical operations like attention and state space models (SSMs), newer ecosystems like Apple Silicon lack this depth of optimized kernels, often leaving significant performance on the table despite the capabilities of the MLX framework.

K-Search: Evolutionary Kernel Optimization

K-Search is an evolutionary framework that uses an iterative loop to optimize GPU kernels. The process involves three primary phases:

  1. Action Selection: An LLM (specifically Gemini 3.5 Pro Preview in this study) acts as a "GPU kernel performance engineer," analyzing the kernel's classification, data layout, and likely bottlenecks to propose an optimization action from a search tree (the "world model").
  2. Local Refinement: A code-writing model generates candidate implementations based on the selected action, which are then compiled and benchmarked on real hardware.
  3. World Model Update: The LLM reasons over the results to update the search tree by inserting new actions, updating priority scores, or pruning unsuccessful paths.

This search is grounded by a "Spec"—a domain-specific document containing hardware rules and mathematical constraints to prevent the generation of invalid primitives.

The CUDA-to-MLX Translation Layer

To bridge the gap between NVIDIA and Apple Silicon architectures, the researchers developed a translation layer that converts CUDA conceptual knowledge into MLX/Metal strategies. This layer consists of:

  • Concept Mapping Tables: A glossary mapping CUDA primitives to Metal equivalents with hardware-specific constraints (e.g., mapping __shared__ memory to Metal threadgroup memory while accounting for the 32 KB limit on Apple Silicon versus 48 KB on NVIDIA).
  • MLX-Specific Hints: Guidance on patterns without direct CUDA equivalents, such as using simd_shuffle_xor for register-based row reductions or the "exp2 trick" (replacing $e^x$ with $2^{x \log_2 e}$) to leverage Apple's fast fast::exp2() hardware instruction.
  • Reusable Assertions: Reframing expert kernel behaviors as properties that the evolutionary search must maintain rather than copying code literally.

Performance Benchmarks

Attention Kernel Results

By providing the evolutionary search with full context from the translation layer, the researchers achieved near-expert performance. The evolved kernel independently discovered and implemented advanced strategies including:

  • Threadgroup memory tiling
  • Online softmax
  • K-transposition for memory access
  • The exp2 trick

This resulted in a performance jump from 0.26x (pure evolution) to 0.97x the speed of Apple's state-of-the-art native attention kernel.

Mamba SSM Kernel Results

K-Search was applied to the Mamba state-space model (SSM) kernel to test generalization. On an M1 Max (64GB) using mamba-370m f16, the evolved mlx-mamba kernel demonstrated a massive increase in prefill throughput compared to the community mlx-lm implementation:

Metric mlx-mamba (ours) mlx-lm (community) mamba.py
Decode 152 tok/s 116 tok/s 40 tok/s
Prefill L=512 5,751 tok/s 329 tok/s 1,089 tok/s
Prefill L=1024 6,010 tok/s 327 tok/s 1,127 tok/s
Prefill L=2048 6,612 tok/s 1,092 tok/s 1,092 tok/s
Prefill L=4096 6,743 tok/s 339 tok/s 1,042 tok/s

Key Insight: The ~20x prefill speedup is attributed to the implementation of a parallel (prefix) scan. While the community mlx-lm implementation processes tokens sequentially, the evolved kernel uses an associative combine to evaluate the sequence in $O(\log N)$ dependent steps, fully utilizing Apple Silicon's GPU throughput during the prefill phase.

Future Directions

The researchers are extending this work to support additional hardware architectures, including the IBM Spyre AIU, and developing more complex kernels such as fused MoE routing and paged attention. The primary finding is that the bottleneck in AI-driven kernel generation is not the LLM's coding ability, but the quality of the architectural context and constraints provided to the model.

Sources