MegaOps lets separate GPU kernels overlap through data tiles
MegaOps: A Kernel Overlapping Framework with PyTorch-like Modularity and Megakernel Performance · arXiv:2611.16248
In plain words
MegaOps is a framework for running the pieces of a machine-learning model as separate GPU kernels. Usually, a later kernel waits for the whole earlier one to finish. MegaOps lets it begin once the particular pieces of shared data it needs are ready.
In tests of short-context decoding on one H100 GPU, MegaOps processed requests at a higher average rate than vLLM across the tested models and batch sizes. The result applies to those workloads, and some separate small-shape tests favored another system.
The research
Model software can be easier to change and debug when its operators remain separate, but whole-kernel waits can slow execution. MegaOps’s results change the expectation that keeping kernels separate necessarily means giving up the performance of fused kernels on the tested serving workloads.
MegaOps lets dependent model kernels overlap by synchronizing on shared data tiles, testing whether separate kernels can keep pace with fused ones.
MegaOps launches dependent kernels separately and has them signal when shared data tiles are ready, so a consumer need not wait for the producer’s entire run. The authors compared its serving throughput with vLLM and specialized fused kernels.
MegaOps averaged higher throughput than vLLM in tested single-GPU decoding and multi-GPU prefill, and reached the reported range of hand-tuned ThunderKittens kernels on their tuned workloads. These results support keeping kernels separate while approaching fused-kernel performance in those settings.
The results cover the tested dense Llama and Qwen models and GPU setups. In a separate small-shape matrix-multiplication and reduce-scatter test, MegaOps lagged TritonDistributed; its grouped no-overlap comparison also changed several features together.
What the evidence shows
- Keep kernels separate
MegaOps is a framework that runs model operators as separate GPU kernels and coordinates dependent kernels through the data they share. This addresses a tradeoff: waiting for a whole kernel can delay useful work, while fusing operators into one large kernel can make them harder to replace and debug. MegaOps allows a consumer kernel to start once the particular input tiles it needs are ready.
- Serving results meet the challenge
On short-context decoding on one H100, the average comparison favors MegaOps over vLLM.
Against hand-tuned ThunderKittens kernels, MegaOps approaches their throughput only on the workloads those kernels were tuned for, showing that separate kernels can compete with specialized fused kernels in these tested serving cases.
- A multi-GPU check
The multi-GPU prefill comparison checks whether the approach also performs in a setting where prompt processing is spread across GPUs.
On 8xH100, MegaOps averaged 1.18 times vLLM’s throughput across the tested Qwen3-8B and Llama-3.1-70B cases and batch sizes, though the systems also differed in kernels and parallel execution scheme.
- What the control can show
In that same multi-GPU prefill setting, full MegaOps reached 1.24 times the throughput of a version with multi-streaming disabled and prefetch and shared-memory padding removed. The control indicates a combined difference associated with those changes.
Since several features changed together, it cannot show how much any one feature contributed on its own.
- Small shapes expose a limit
A separate matrix-multiplication and reduce-scatter test checks a workload beyond model serving.
For Llama-3.1-8B at 4K tokens or fewer on 8xH100, MegaOps trailed TritonDistributed, which switches to a different matrix-multiplication kernel; the serving results therefore do not establish an advantage on every workload.
Deep
- Original paper authors
- Jinkun Lin, Byungsoo Oh, Rachee Singh
- Submitted to arXiv
- License
- CC BY 4.0