Optimal Software Pipelining and Warp Specialization for Tensor Core GPUs
Not on the map. Why?
Abstract (arXiv)
GPU architectures have continued to grow in complexity, with recent incarnations introducing increasingly powerful fixed-function units for matrix multiplication and data movement to accompany highly parallel general-purpose cores. To fully leverage these machines, software must use sophisticated schedules that maximally utilize all hardware resources. Since realizing such schedules is complex, both programmers and compilers routinely employ program transformations, such as software pipelining (SWP) and warp specialization (WS), to do so in practice. However, determining how best to use SWP and WS in combination is a challenging problem that is currently handled through a mix of brittle compilation heuristics and fallible human intuition, with little insight into the space of solutions. To remedy this situation, we introduce a novel formulation of SWP and WS as a joint optimization problem that can be solved holistically by off-the-shelf constraint solvers. We reify our approach in Twill, the first system that automatically derives optimal SWP and WS schedules for a large class of iterative programs. Twill is heuristic-free, easily extensible to new GPU architectures, and guaranteed to produce optimal schedules. We show that Twill can rediscover, and thereby prove optimal, the SWP and WS schedules manually developed by experts for Flash Attention on both the NVIDIA Hopper and Blackwell GPU architectures.
Key facts. Cited by 4 papers in the map (data through 7 Oct 2026). Builds on 8 of its 36 references that are in the map. First posted on arXiv 19 Dec 2025.
Builds on
8 of 36 references.
References · 8
- 2019PyTorch: An Imperative Style, High-Performance Deep Learning Librarycited by 20,160
- 2022FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awarenesscited by 3,611
- 2023FlashAttention-2: Faster Attention with Better Parallelism and Work Partitioningcited by 2,551
- 2024ThunderKittens: Simple, Fast, and Adorable AI Kernelscited by 34
- 2025Tawa: Automatic Warp Specialization for Modern GPUs with Asynchronous Referencescited by 6
- 2022ALCOP: Automatic Load-Compute Pipelining in Deep Learning Compiler for AI-GPUscited by 4
- 2025Task-Based Tensor Computations on Modern GPUscited by 4
- 2025The SCIP Optimization Suite 10.0cited by 4
Built on by
Citing papers per year: 2026: 4.
4 papers.
- 2026Magellan: Autonomous Discovery of Novel Compiler Optimization Heuristics with AlphaEvolvecited by 5
- 2026CAKE: Compiler-Agent Co-Design for Frontier Kernel Evolutioncited by 2
- 2026GauS: Differentiable Scheduling Optimization via Gaussian Reparameterizationcited by 0
- 2026Hardware-Aware Features for CUTLASS Kernel Selectioncited by 0
- 2026Hardware-Aware Features for CUTLASS Kernel Selection28 Sep 2026
- 2026CAKE: Compiler-Agent Co-Design for Frontier Kernel Evolution12 Aug 2026
- 2026GauS: Differentiable Scheduling Optimization via Gaussian Reparameterization23 Feb 2026
- 2026Magellan: Autonomous Discovery of Novel Compiler Optimization Heuristics with AlphaEvolve28 Jan 2026