# CODA: Rewriting Transformer Blocks as GEMM-Epilogue Programs

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

## Key facts

- Cited by 5 papers in the map (data through 14 Oct 2026).
- Builds on 18 of its 24 references that are in the map.
- First posted on arXiv 19 May 2026.

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

- [The Llama 3 Herd of Models](https://mlsanity.com/paper/2407.21783/) (2024)
- [Efficient Memory Management for Large Language Model Serving with PagedAttention](https://mlsanity.com/paper/2309.06180/) (2023)
- [RoFormer: Enhanced Transformer with Rotary Position Embedding](https://mlsanity.com/paper/2104.09864/) (2021)
- [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)
- [KernelBench: Can LLMs Write Efficient GPU Kernels?](https://mlsanity.com/paper/2502.10517/) (2025)
- [Liger Kernel: Efficient Triton Kernels for LLM Training](https://mlsanity.com/paper/2410.10989/) (2024)
- [Learning to Discover at Test Time](https://mlsanity.com/paper/2601.16175/) (2026)
- [Data Movement Is All You Need: A Case Study on Optimizing Transformers](https://mlsanity.com/paper/2007.00072/) (2020)
- [TorchTitan: One-stop PyTorch native solution for production ready LLM pre-training](https://mlsanity.com/paper/2410.06511/) (2024)
- [Towards Robust Agentic CUDA Kernel Benchmarking, Verification, and Optimization](https://mlsanity.com/paper/2509.14279/) (2025)
- [TileLang: A Composable Tiled Programming Model for AI Systems](https://mlsanity.com/paper/2504.17577/) (2025)
- [Mirage: A Multi-Level Superoptimizer for Tensor Programs](https://mlsanity.com/paper/2405.05751/) (2024)
- [Cut Your Losses in Large-Vocabulary Language Models](https://mlsanity.com/paper/2411.09009/) (2024)
- [ThunderKittens: Simple, Fast, and Adorable AI Kernels](https://mlsanity.com/paper/2410.20399/) (2024)
- [CUDA-LLM: LLMs Can Write Efficient CUDA Kernels](https://mlsanity.com/paper/2506.09092/) (2025)
- [CUDA-L2: Surpassing cuBLAS Performance for Matrix Multiplication through Reinforcement Learning](https://mlsanity.com/paper/2512.02551/) (2025)
- [ParallelKittens: Systematic and Practical Simplification of Multi-GPU AI Kernels](https://mlsanity.com/paper/2511.13940/) (2025)

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

- [Sol Video Inference Engine: Agent-Native Full-Stack Acceleration Framework for Efficient Video Generation](https://mlsanity.com/paper/2606.23743/) (2026)
- [Full-Stack FP4: Stable LLM Pretraining with Quantized Projections, Optimizers, and Attention](https://mlsanity.com/paper/2607.04422/) (2026)
- [FORGE: Fused On-Register Gradient Elimination for Memory-Efficient LLM Training](https://mlsanity.com/paper/2606.22932/) (2026)
- [Sol-H3: Recursive Self-Improvement for MiniMax-H3 Inference Acceleration on Sol-Engine across Cloud and Edge](https://mlsanity.com/paper/2609.35110/) (2026)
- [Format-Aware Fusion for Fast FP4 Pretraining](https://mlsanity.com/paper/2610.00053/) (2026)

## Abstract (arXiv)

Transformer training systems are built around dense linear algebra, yet a nontrivial fraction of end-to-end time is spent on surrounding memory-bound operators. Normalization, activations, residual updates, reductions, and related computations repeatedly move large intermediate tensors through global memory while performing little arithmetic, making data movement an increasingly important bottleneck in otherwise highly optimized training stacks. We introduce CODA, a GPU kernel abstraction that expresses these computations as GEMM-plus-epilogue programs. CODA is based on the observation that many Transformer operators exposed as separate framework kernels can be algebraically reparameterized to execute while a GEMM output tile remains on chip, before it is written to memory. The abstraction fixes the GEMM mainloop and exposes a small set of composable epilogue primitives for scaling, reductions, pairwise transformations, and accumulation. This constrained interface preserves the performance structure of expert-written GEMMs while remaining expressive enough to cover nearly all non-attention computation in the forward and backward pass of a standard Transformer block. Across representative Transformer workloads, both human- and LLM-authored CODA kernels achieve high performance, suggesting that GEMM-plus-epilogue programming offers a practical path toward combining framework-level productivity with hardware-level efficiency.