PipeLink enables flexible linking of GPU kernel parts without performance loss
PipeLink: Contract-Guided Linking of GPU Kernel Components
Distributed, Parallel, and Cluster Computing
Summary
When running complex GPU programs for machine learning, different parts like loaders and math operations usually have to be combined all at once, making it hard to swap or share parts. The authors created PipeLink, which lets these parts be compiled separately and then safely linked together by checking their requirements and fitting them onto the GPU. This makes it easier to build new combinations of GPU code without losing speed. Their tests show PipeLink performs almost as well as compiling everything as one piece, and can improve performance for certain tasks.
What this means in practice
- •For gpu software engineers: Build and maintain GPU machine-learning kernels by linking modular components without needing full recompilation or complex manual code merging.
- •For machine learning infrastructure teams: Combine optimized attention and decoding kernel components from different sources to create faster and flexible training and inference pipelines.
Authors
Jize Guo, Xupeng Miao
Abstract
The fastest machine-learning kernels are compositions of components, such as loaders, matrix-multiply mainloops, epilogues and communication stages, that run concurrently on one GPU streaming multiprocessor and share its on-chip memory, barriers, warps and registers. Because these components hard-code where they sit in this shared space, they cannot be compiled, shipped or replaced on their own: combining components from different sources means forking them by hand or keeping them in separate kernels. We present PipeLink, a contract-guided linker for GPU kernel components. Each component is compiled separately into an LTO object that carries a contract stating what it requires and what it promises; PipeLink checks the contracts against each other, relocates every component by solving one joint placement of the SM's resources, and generates the transfers between them, and the vendor LTO backend folds the result into constants. Across 51 configurations on Hopper and Blackwell GPUs, linked kernels run within 0.2% of the same sources compiled as one unit. Linking makes new compositions practical: a CuTe attention operator inside a ThunderKittens decode megakernel speeds up decoding by 2.16$\times$, a four-owner MoE training megakernel matches the original monolith, one unchanged attention core yields sparse and context-parallel variants that outperform FlexAttention, vLLM and Transformer Engine, and a GEMM + SwiGLU composition whose consumer serves three GPU architectures runs 1.2$\times$ faster on Blackwell than cuBLAS with a compiled activation.