# MPK: A Compiler and Runtime for Mega-Kernelizing Tensor Programs

arXiv: https://arxiv.org/abs/2512.22219 · HTML: https://mlsanity.com/paper/2512.22219/

## Key facts

- Cited by 15 papers in the map (data through 14 Oct 2026).
- Builds on 14 of its 47 references that are in the map.
- First posted on arXiv 22 Dec 2025.

## Builds on (most cited first, 14 of 14)

- [LLaMA: Open and Efficient Foundation Language Models](https://mlsanity.com/paper/2302.13971/) (2023)
- [Megatron-LM: Training Multi-Billion Parameter Language Models Using Model Parallelism](https://mlsanity.com/paper/1909.08053/) (2019)
- [SGLang: Efficient Execution of Structured Language Model Programs](https://mlsanity.com/paper/2312.07104/) (2023)
- [FlashInfer: Efficient and Customizable Attention Engine for LLM Inference Serving](https://mlsanity.com/paper/2501.01005/) (2025)
- [Ansor: Generating High-Performance Tensor Programs for Deep Learning](https://mlsanity.com/paper/2006.06762/) (2020)
- [Learning to Optimize Tensor Programs](https://mlsanity.com/paper/1805.08166/) (2018)
- [Mirage: A Multi-Level Superoptimizer for Tensor Programs](https://mlsanity.com/paper/2405.05751/) (2024)
- [TensorIR: An Abstraction for Automatic Tensorized Program Optimization](https://mlsanity.com/paper/2207.04296/) (2022)
- [ThunderKittens: Simple, Fast, and Adorable AI Kernels](https://mlsanity.com/paper/2410.20399/) (2024)
- [Equality Saturation for Tensor Graph Superoptimization](https://mlsanity.com/paper/2101.01332/) (2021)
- [Triton-distributed: Programming Overlapping Kernels on Distributed AI Systems with the Triton Compiler](https://mlsanity.com/paper/2504.19442/) (2025)
- [TokenWeave: Efficient Compute-Communication Overlap for Distributed LLM Inference](https://mlsanity.com/paper/2505.11329/) (2025)
- [ParallelKittens: Systematic and Practical Simplification of Multi-GPU AI Kernels](https://mlsanity.com/paper/2511.13940/) (2025)
- [Optimal Kernel Orchestration for Tensor Programs with Korch](https://mlsanity.com/paper/2406.09465/) (2024)

## Built on by (most cited first, 15 of 15)

- [Event Tensor: A Unified Abstraction for Compiling Dynamic Megakernel](https://mlsanity.com/paper/2604.13327/) (2026)
- [Ada-MK: Adaptive MegaKernel Optimization via Automated DAG-based Search for LLM Inference](https://mlsanity.com/paper/2605.11581/) (2026)
- [Vortex: Efficient and Programmable Sparse Attention Serving for AI Agents](https://mlsanity.com/paper/2606.06453/) (2026)
- [ExpertPlex: A High-Goodput Disaggregated Serving System for MoE LLMs with Adaptive Persistent Kernels](https://mlsanity.com/paper/2607.18002/) (2026)
- [PhyAI: Real-Time Physical AI at the Edge, Scalable Rollouts in the Cloud](https://mlsanity.com/paper/2608.03682/) (2026)
- [Ave: Guiding Agentic GPU Optimization Using Data-Flow Invariants](https://mlsanity.com/paper/2604.18616/) (2026)
- [AutoMegaKernel: A Statically-Checked Agent Harness for Self-Retargeting Megakernel Synthesis](https://mlsanity.com/paper/2606.09682/) (2026)
- [Measuring and Reducing WebGPU Dispatch Overhead for LLM Inference](https://mlsanity.com/paper/2608.08730/) (2026)
- [Ready Cohorts: Bounding GPU Opportunity and Avoiding Host Round Trips in LLM-Agent Control](https://mlsanity.com/paper/2608.12123/) (2026)
- [mKernel: Fast Multi-GPU, Multi-Node Fused Kernels](https://mlsanity.com/paper/2609.13585/) (2026)
- [Empowering Hybrid Attention Models on NPUs](https://mlsanity.com/paper/2609.32114/) (2026)
- [Mixture-of-Kittens: MoE Megakernel for NVL72s](https://mlsanity.com/paper/2609.36070/) (2026)
- [Purlin: Separating Orchestration from the Datapath of Collectives](https://mlsanity.com/paper/2609.36954/) (2026)
- [Trident: Unifying Guarded Dispatch and Host Execution for PyTorch Triton Workloads](https://mlsanity.com/paper/2609.37241/) (2026)
- [RESOLVE: Language-Agnostic Validation of GPU Kernels Through Testing, Reduction, and Proof](https://mlsanity.com/paper/2610.05683/) (2026)

## Abstract (arXiv)

We introduce Mirage Persistent Kernel (MPK), the first compiler and runtime system that automatically transforms multi-GPU model inference into a single high-performance mega-kernel. MPK introduces an SM-level graph representation that captures data dependencies at the granularity of individual streaming multiprocessors (SMs), enabling cross-operator software pipelining, revfine-grained overlap of computation and communication, and other optimizations that are infeasible under the conventional kernel-per-operator execution model. The MPK compiler lowers tensor programs into optimized SM-level task graphs and generates fast CUDA implementations for each task, while the MPK in-kernel parallel runtime executes these tasks within a single persistent mega-kernel using decentralized scheduling across SMs. Together, these components provide end-to-end kernel fusion with minimal developer effort, while preserving the flexibility of existing programming models. Our evaluation shows that MPK significantly outperforms existing kernel-per-operator LLM serving systems, achieving up to 1.7× lower end-to-end inference latency and pushing LLM inference performance close to the limits of the underlying hardware. MPK is publicly available at https://github.com/mirage-project/mirage.