Understanding GPU Megakernels: Why Residency Can Mislead

Abstract

Megakernels let one persistent GPU program prepare dependent operators early and execute ready tasks together. We study the benefits and costs of these capabilities by reconstructing two LLM megakernels, ThunderKittens (TK) and VDCores, as separate kernels on H100 using their original operator code. Our results show that CTA residency can misrepresent opportunities for useful overlap in both directions. By overlapping startup and operand loading with programmatic dependent launch (PDL), the reconstructions reduce decode-step latency by 5.7% in TK and achieve comparable performance in VDCores when both executions enforce the required memory ordering. Under whole-operator dependencies, TK's persistent execution still gains more from prefetch, but runtime and completion-signal ordering costs outweigh this advantage. Both startup overlap and prefetch depend on resource availability. Persistent CTAs can reuse storage before retiring, whereas separate successors must acquire their full CTA footprints before initializing or loading operands. This admission constraint in both reconstructions motivates finer resource sharing across kernels. Beyond early preparation, joint execution requires ready work and spare capacity in the bottleneck resource. Reserving SMs for communication can add GEMM waves that erase overlap gains, while filling idle SMs offers little benefit once memory bandwidth is saturated. These findings clarify benefits and costs of megakernels and motivate GPU support for earlier resource reuse across kernels.