PipeLink: Contract-Guided Linking of GPU Kernel Components

• Components hard-code shared-memory ownership, barriers, warps, registers, and tile assumptions, so cross-source compositions require hand-forking or separate kernels.Separate-kernel context-parallel attention also incurs launch, memory-traffic, and lost-overlap costs at each boundary.Source: §1 · p. 1
• Each separately compiled LTO object carries a contract; PipeLink checks contracts, jointly places components, and generates transfers between them.The vendor LTO backend folds the placement into constants, making linking cost nothing at run time.Source: §1 · p. 2
• Across 51 timed configurations, linked kernels run within 0.2% of the same sources compiled as a single unit.The evaluation spans Hopper and two generations of Blackwell GPUs.Source: §1 · p. 2
• A CuTe attention operator decodes 2.16× faster at 4096-token context, while the four-owner MoE megakernel matches its monolith at 0.99–1.01×.Source: §1 · p. 2
• One unchanged attention core yields sparse attention up to 38× faster than FlexAttention and context-parallel attention 1.20–1.31× faster than Transformer Engine.Source: §1 · p. 2
• The Qwen3-32B GEMM + SwiGLU composition runs 1.20× faster on B300 than cuBLAS followed by a compiled SwiGLU.Its consumer links with producers on three GPU architectures.Source: §1 · p. 2

2.1 Anatomy of a Fused Kernel

• A context-parallel attention kernel runs loader, attention-compute, epilogue and communication warps concurrently, coordinating through shared on-chip resources and synchronization.The running example uses twelve warps and shares SM resources across components.Source: §2.1 · pp. 2–3

2.2 Components Are Absolute-Addressed

• Shared memory, barriers, warps and registers are assigned jointly in fused kernels, so independently written components collide when each assumes its own starting positions.The same assignments are hard-coded into shared-storage layouts, barrier IDs, warp roles and register budgets.Source: §2.2 · p. 3, §2.1 · p. 3

2.3 Fork or Split

• Forking ports an outside operator into one kernel’s conventions and creates another kernel to maintain; splitting preserves components but pays for launches, memory round trips and lost overlap.Production context-parallel attention uses separate kernels and operations for key/value blocks, communication and output correction.Source: §2.3 · p. 3

2.4 Symbol Linking Is Not Enough

• Three of five symbol-linked MoE combinations hung, while two completed with unwritten output after one owner changed a geometry constant.The four owners schedule against shared tile geometry, which symbol linking does not record.Source: §2.4 · p. 3, Figure 2 · p. 3

3.1 Components

• A component is a separately compiled fused-kernel part, including a warp-group role, a callable, or a tensor producer or consumer.Components may be written in CUDA, CUTLASS/CuTe, ThunderKittens or TileLang and distributed without source in an LTO object.Source: §3.1 · p. 4

3.2 Contract Vocabulary

• Contracts describe resource claims by handle, lifetimes, synchronization protocols and typed constants shared between owners and consumers.Claims specify needs rather than addresses; typed facts let the linker propagate an owner’s changed tile size to consumers without recompiling them.Source: §3.2 · p. 4, Table 1 · p. 4

3.3 Relocatable Components

• Handle functions let components read placements chosen per composition; the LTO backend replaces their link-time bindings with constants, avoiding runtime loads.This keeps objects unchanged across compositions while making placement values compile to immediates.Source: §3.3 · p. 4

3.4 Authoring

• Authors annotate only semantic promises the compiler cannot observe, while packing captures measurable object facts and rejects layout maps that do not fit exactly.TileLang can infer tensor-import demand from the destination layout; in the SwiGLU example, its consumer source is 56 lines versus 163 in CUDA.Source: §3.4 · p. 5
• The four MoE megakernel owners use 327 annotation statements across their components and host provider, against 3,482 lines of source.Other case-study components use 11–20 annotation statements.Source: §3.4 · p. 5

4.1 Checking

• The linker checks paired endpoints for matching memory space, scope, tile, layout and protocol, and rejects conflicts with a clause and the conflicting facts.It also checks callable signatures, owner-consumer facts and compatible topologies; uncalled exports are dropped.Source: §4.1 · p. 5, Figure 3 · p. 5

4.2 Relocation as Constraint Solving

• The solver jointly places resource claims subject to capacity, alignment, warp coverage and non-overlap, allowing reuse when phases or tensor lifetimes do not overlap.It prefers deeper rings and larger register budgets, then a smaller footprint, and produces deterministic placements.Source: §4.2 · pp. 5–6
• First-fit packing misses the smallest MoE shared-memory footprint; the solver finds it, leaving 1,440 of 232,448 bytes free.Across compositions with 4–42 handles, the solver proves objectives optimal within 1.3% of its work limit.Source: §4.2 · p. 6

4.3 Why Linking Costs Nothing

• Handle stubs, call trampolines and a dispatch kernel are LTO-linked with components, letting the backend fold solved placements into constants.The linker does not reschedule components; if backend register allocation does not fit, it re-solves once and otherwise fails with L-REGS.Source: §4.3 · p. 6, §4.2 · p. 6

4.4 Generated Tensor Transfers

• The linker maps producer exports to consumer demands and lowers the first verified route, from direct register or tensor-memory access to shuffle or shared-memory staging.Candidate routes are checked for exact delivery, race freedom and ordering; generated transfers do not convert values.Source: §4.4 · p. 6, Figure 4 · p. 7

4.5 Communication and Host Interface

• Communication components declare the domains, allocation views and readiness counters they use; the linker checks agreement and binds caller-provided addresses without synthesizing collectives.Host providers fill their own parameter regions from caller-lent buffers, and an opened session can be captured in a CUDA graph.Source: §4.5 · pp. 6–7

5.1 Setup

• The evaluation compares each linked kernel with a single-unit control built from the same component sources, generated connecting code, and placement; linked kernels match controls and pass memcheck, racecheck, and synccheck.The compiler uses nvcc, Z3, TileLang, and nvJitLink; experiments lock each GPU to its specified sustained SM clock.Source: §5.1 · p. 7
• Across 51 timed configurations, linked kernels run within 0.2% of their single-unit controls, and normalized machine code is identical for the listed attention, megakernel, and GEMM–SwiGLU compositions.Experiments use H200, B300, and RTX 5090 GPUs; linked kernels match controls bitwise or within megakernel run-to-run noise and pass the reported sanitizer checks.Source: §5.1 · p. 7

5.2 Cross-Library Megakernel

• At a 4096-token context, the mixed-library decode kernel is 2.16× faster than the ThunderKittens kernel; at 16 tokens, it is 5% slower.The CuTe attention package replaces one of nine packages while the other packages, instruction stream, and host provider remain unchanged; at 1024 tokens with split-KV reduction it is 1.22× faster.Source: §5.2 · pp. 7–8, Table 3 · p. 8

5.3 Multi-Owner MoE Megakernel

• On four B300 GPUs, 12 MoE kernels reach 0.989–1.014× the throughput of the original monolithic step, while matching single-unit controls at 0.999–1.001×.Six packages link into 24 kernels; an independent SwiGLU owner also links unchanged with the other packages and passes all 48 per-rank checks, while incompatible combinations are rejected with named diagnostics.Source: §5.3 · p. 8

5.4 Attention Variants from One Core

• Sparse attention reaches up to 38× the speed of FlexAttention, while context-parallel attention is 1.20–1.31× faster than Transformer Engine at 8K and 32K tokens.The seven variants reuse one unchanged attention core; MSA is 1.6–4.5× faster than vLLM’s dedicated kernel, and context-parallel attention is 1.95× faster on two H200 GPUs and 3.67× on four than the single-GPU kernel for 32K-token causal prefill.Source: §5.4 · pp. 8–9, Table 4 · p. 9, Table 5 · p. 9

5.5 Fused GEMM and Activation

• On B300, the GEMM–SwiGLU composition is 1.20–1.21× faster than cuBLAS followed by compiled SwiGLU and 1.34–1.37× faster than cuBLAS followed by Liger’s SwiGLU.Against hand-fusing the same components into a native kernel, linked kernels reach 0.94–1.01× its throughput; all 12 configurations are within 0.9987–1.0015× of single-unit controls.Source: §5.5 · p. 9, Table 6 · p. 9

In one source program

• Template libraries, DSLs, and fusion systems compose within one program, while PipeLink links separately compiled components that may use those same tools.Existing extension points are designed by the library authors; PipeLink works across independently compiled components.Source: §6 · pp. 9–10

In compiler IR

• Warp-specialization and layout compilers own or reschedule the whole kernel; PipeLink preserves each opaque component’s compiled schedule and uses declared interfaces.Those compilers partition roles or represent layouts connecting operations, whereas PipeLink links components compiled separately.Source: §6 · p. 10

In runtime task graphs

• Runtime task-graph systems decide which task runs where, while PipeLink decides how separately compiled components share one SM.Their task bodies could be linked components, as in the Llama-1B megakernel.Source: §6 · p. 10

By symbol and communication overlap

• Device LTO links separately compiled functions by symbol, whereas PipeLink adds contracts for shared resources and checks and places components; it adds no communication algorithm.PipeLink can let one owner’s communication component join another owner’s compute components in one kernel.Source: §6 · p. 10

Synchronization verification

• Synchronization-verification tools check whole-program synchronization or schedule safety; PipeLink checks declared inter-component protocols and leaves code checking to sanitizers and numerical tests.Source: §6 · p. 10

7 LIMITATIONS

• PipeLink checks contracts against each other, not against component code, so a component that violates its contract can still produce an incorrect kernel.Numerical checks and sanitizers remain part of each composition, as with hand-fused kernels.Source: §7 · p. 10
• Generated code only moves data within one thread block; arithmetic, precision conversion, and cross-block or cross-GPU communication remain in component code.Tensor exports carry one generation, so a producer handing off its accumulator this way cannot be persistent; multi-slot rings must be written in component code.Source: §7 · p. 10

7 LIMITATIONS AND CONCLUSION

• PipeLink preserves component schedules rather than making components faster, and packages are compiled for one architecture and toolkit because LTO IR is not portable across CUDA major versions.Existing kernels also need to be ported to declared interfaces.Source: §7 · p. 10