PipeLink: Contract-Guided Linking of GPU Kernel Components

P IPE L INK : C ONTRACT-G UIDED L INKING OF GPU K ERNEL C OMPONENTS

arXiv:2610.37200v1 [cs.DC] 27 Oct 2026

Jize Guo 1 Xupeng Miao 1

A BSTRACT 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×, 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× faster on Blackwell than cuBLAS with a compiled activation. 1

I NTRODUCTION

The fastest machine-learning kernels on modern GPUs are no longer single operators. An attention kernel in the FlashAttention family runs a tile loader, a matrix-multiply mainloop, a softmax stage and an epilogue as groups of warps that execute concurrently and hand tiles to each other through on-chip memory (Shah et al., 2024; Zadouri et al., 2026). Megakernels go further and run an entire decode step or training pass inside one persistent kernel (Spector et al., 2025a; Cheng et al., 2026; Jin et al., 2026; Sul et al., 2026b), and communication-overlapped kernels move data between GPUs while they compute (Chang et al., 2024; Zheng et al., 2025b; Sul et al., 2026a). Each such kernel is a composition of components that cooperate at fine grain on the same streaming multiprocessor (SM). Yet these components cannot be built, shipped or replaced as components. They agree on physical facts that no interface records: which bytes of shared memory each one owns, which barriers and warps it uses, how many registers it may hold, and which tile sizes its neighbors schedule against. Whoever wants to combine components from dif1 School of Computer Science, Peking University, Beijing, China. Correspondence to: Xupeng Miao <xupeng.miao@pku.edu.cn>.

Preprint. Under review.

ferent sources therefore faces a dilemma: fork the sources and re-derive these facts by hand, or keep the components in separate kernels. Megakernels take the first path: they are written as one translation unit, by one team, in one library, so adopting a faster operator from another library, such as a CuTe (Thakkar et al., 2023) attention in a ThunderKittens (Spector et al., 2025b) megakernel (Spector et al., 2025a), means porting it to the megakernel’s conventions by hand. New attention variants also start as forks of an expert kernel, and systems that move context-parallel communication into the kernel write or modify their own attention kernel (Sul et al., 2026a; Qiang et al., 2026). Production context-parallel attention takes the second path: it calls an unmodified FlashAttention kernel once per key/value block and runs the transfers between GPUs and the correction of partial outputs as separate operations around it (NVIDIA, f; zhuzilin; Fang & Zhao, 2024), paying for every kernel boundary in launches, memory traffic and lost overlap. Existing tools compose kernels either before or below this boundary. Template libraries, domain-specific languages and warp-specialization compilers compose operations within one source program that they own (Thakkar et al., 2023; Dong et al., 2025; Wang et al., 2026; Chen et al., 2026b; Soi et al., 2026; Yadav et al., 2025). Device linking with link-time optimization (LTO) does combine separately compiled code into one kernel at no runtime cost (Sundaram et al., 2023), but it links by symbol: it records which names

PipeLink: Contract-Guided Linking of GPU Kernel Components Compiled components

Linked kernels Attention, 1 GPU core smem base() = 0 smem

load

core

epilogue

PipeLink

warps

L

check

KV push

CP load

attention team load author communication team ring counter

rings core

E

same compiled core

relocate generate

Context-parallel attention, GPU 0 of 4 core smem base() = 1024 smem rings warps

L

E

push

core

GPU 1 GPU 2 GPU 3

NVLink

Figure 1. PipeLink overview. Separately compiled components carry contracts at their interfaces (left). PipeLink checks the contracts, relocates the components and generates the code between them, so the same compiled attention core lands at different shared-memory offsets in two kernels (right). Lane widths are schematic.

each fragment defines and needs, and nothing about the resources it occupies or the facts its neighbors compiled against. When one owner of a mixture-of-experts (MoE) training kernel changes a tile size and recompiles alone, device linking accepts the stale combination, and the kernel hangs or silently leaves outputs unwritten (Section 2.4). We argue that kernel components need what object files gave whole programs: a contract that states what crosses the component boundary, and a linker that relocates each component into the kernel it joins. Today’s components are, in effect, absolute-addressed: they hard-code their positions in a small on-chip address space that every other component shares. A relocatable component instead declares what it requires and what it promises, and reads its placement through bindings fixed only at link time. Our thesis is that such contracts, checked and bound at link time, let separately compiled components share one SM as efficiently as if their sources had been merged by hand. This paper presents PipeLink, a contract-guided linker for GPU kernel components. Each component is compiled separately into an LTO object that carries its contract. PipeLink checks the contracts against each other, solves one joint placement of all components’ on-chip resources for the composition at hand, generates the code that moves tensors between components, and binds the memory that communication components share across GPUs. The vendor LTO backend then folds the placement into constants, so linking costs nothing at run time. We evaluate PipeLink on Hopper and two generations of Blackwell GPUs. Across 51 timed configurations, linked kernels run within 0.2% of the same sources compiled as a single unit, and compositions that previously required a fork become practical: • a CuTe attention operator inside a ThunderKittens Llama-1B decode megakernel decodes 2.16× faster at a 4096-token context;

• an MoE training megakernel split among four owners matches the original monolith (0.99–1.01×) and survives every stale-owner upgrade; • 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; • a fused GEMM and SwiGLU for Qwen3-32B, whose consumer links with producers on three GPU architectures, runs 1.20× faster on B300 than cuBLAS followed by a compiled SwiGLU. In summary, this paper makes the following contributions: • Component contracts, which state what a concurrently running GPU kernel component requires and promises, so that it can be compiled, distributed and replaced on its own, without its source (Section 3). • A contract-guided linker that checks contracts against each other, relocates components by jointly solving their placement, and generates the data transfers between them (Section 4). • Zero-cost composition of real ML kernels across libraries, owners and GPU architectures (Section 5).

2

M OTIVATION

2.1

Anatomy of a Fused Kernel

Our running example is context-parallel attention on Hopper (Figure 1). Its core is a warp-specialized attention kernel derived from FlashAttention-3 (Shah et al., 2024): a load warp copies query, key and value tiles into shared memory, two consumer warpgroups compute attention with an online softmax, and an epilogue warp writes the output. In the context-parallel kernel, two more warps push the local key/value shard to the other GPUs. All twelve warps run

PipeLink: Contract-Guided Linking of GPU Kernel Components

concurrently and hand work to each other through mbarrier rings, named barriers and, across GPUs, readiness counters. They share the SM’s shared memory, barriers, warps and register file, and how much of each a component may use depends on which other components share the thread block.

framework DISPATCH Nb COMBINE Nb

communication DISPATCH Nb

activation

GEMM arrive

arrive

2.2

PIPE DEPTH

routed buffers done

arrive

Components Are Absolute-Addressed Upgrade: communication repacks alone, DISPATCH Nb 512 → 256

Fused-kernel code today fixes all of these assignments in its source. A FlashAttention (Shah et al., 2024) or CUTLASS (Thakkar et al., 2023) kernel declares one SharedStorage struct that lays out every buffer and barrier at a constant offset, numbers its named barriers in an enumeration, selects roles by comparing the warp index against constants, and calls setmaxnreg with constant budgets. ThunderKittens (Spector et al., 2025b) megakernels (Spector et al., 2025a) likewise compile a fixed page and semaphore layout into every operator. In the vocabulary of linkers, these components are absolute-addressed: they hard-code their positions in an address space that is small, shared and spread over several resource kinds. Absolute addressing is harmless as long as one team owns the whole kernel. It becomes the obstacle as soon as a component is meant to be reused. Components written independently each start at shared-memory offset 0, barrier ID 1 and warp 0, so combining them unchanged collides by construction; someone must re-derive every position with knowledge of all the others. 2.3

Fork or Split

Because components cannot be relocated, combining components from different sources leaves two routes: fork their sources into one kernel, or keep them in separate kernels. Fork. Forking merges the sources into one translation unit and re-derives every shared assignment by hand. Megakernels take this route: one team owns an entire model step (Spector et al., 2025a; Cheng et al., 2026; Jin et al., 2026; Sul et al., 2026b; Zheng et al., 2026), and adopting an operator from another library means porting it to the megakernel’s page, semaphore and instruction conventions; recent work automates the port with coding agents that write a new megakernel per model (Li et al., 2026a; Jaber & Jaber, 2026). New attention variants likewise start as forks of an expert kernel such as FlashAttention (Shah et al., 2024; Zadouri et al., 2026), and systems that move context-parallel communication into the kernel write or modify their own attention kernel (Sul et al., 2026a; Qiang et al., 2026; Hu et al., 2026). Every fork is a new kernel to maintain. Split. Splitting keeps each component unmodified in its own kernel and connects the kernels through global memory. Production context-parallel attention takes this route: it calls

linked by symbol symbols only

hangs framework 512

comm 256

linked by contract fact propagated

passes framework 256

comm 256 GEMM 512

GEMM 256

Figure 2. Facts that symbol linking does not record, in the fourowner MoE megakernel. Top: constants that one owner compiles and another schedules against, and a counter shared by three owners. Bottom: the communication owner halves its dispatch tile; linked by symbol the kernel hangs, linked by contract it passes.

an unmodified FlashAttention kernel once per key/value block, moves blocks between GPUs with collective operations on other streams, and corrects the partial outputs and log-sum-exp values in separate operations (NVIDIA, f; zhuzilin; Fang & Zhao, 2024; Yang et al., 2025). Splitting preserves reuse but pays for every boundary with launches, memory round trips and lost overlap. Template libraries avoid both routes only for the variants their authors anticipated, such as masks and score functions over a generic attention kernel (Dong et al., 2025; Ye et al., 2025). 2.4

Symbol Linking Is Not Enough

CUDA already supports separate compilation: relocatable device code and link-time optimization (LTO) combine independently compiled device functions into one kernel as fast as a single translation unit (Scudiero & Murphy, 2014; Sundaram et al., 2023), and NVIDIA’s own libraries ship LTO fragments for this purpose (NVIDIA, c;b). But device linking resolves symbols: it records which names a fragment defines and needs, and nothing about the warps, shared memory, barriers or registers it occupies, or about the facts its neighbors compiled against. The consequence is silent failure. The four owners of our MoE megakernel (Section 5.3) schedule against each other’s tile geometry (Figure 2). We changed one geometry constant of one owner at a time, repacked only that owner, and linked it with the others’ unchanged objects. Device linking accepted all five combinations: three kernels hang, and two complete without error but leave part of their output unwritten. Separately compiled components need an interface that carries what they share, and a linker that acts on it.

PipeLink: Contract-Guided Linking of GPU Kernel Components

3

C OMPONENT C ONTRACTS

3.1

Components

A component is a separately compiled part of a fused kernel. PipeLink supports three kinds. A role is a group of warps with its own entry function, such as a TMA loader, an MMA mainloop or a communication stage. A callable is a device function that runs on its caller’s threads, such as an attention mask, a score modifier or an operator body inside a megakernel interpreter. A tensor producer or consumer is a role or callable that exports or imports a tile held in on-chip storage. One source file may define several of these, and a component may be written in plain CUDA, CUTLASS/CuTe, ThunderKittens or TileLang. The author packs each component into a package: an LTO object produced by the vendor compiler, plus a manifest that records its contract. Packing compiles the source twice, once to evaluate the annotations that make up the contract and once to produce the object. PipeLink never reads or rewrites the component’s instructions; to the linker the object is opaque, and the component can be distributed without its source. The vendor LTO backend, which does see the object’s IR, later optimizes the composed kernel as a whole. We considered composing at two other levels. Composing source or compiler IR, as template libraries and warpspecialization compilers do, gives the composer full freedom but requires every component’s source in one compilation, under one toolchain and one owner. Rewriting compiled machine code would avoid that, but the GPU instruction set is undocumented and changes between generations, and patched binaries miss the cross-component optimizations that make a fused kernel fast. LTO objects sit between the two: they are a supported, stable vendor format, they hide the source, and the backend still optimizes across component boundaries. 3.2

Contract Vocabulary

A contract states what a component requires from the kernel it joins and what it promises to the other components (Table 1). Four properties of this vocabulary matter. First, every resource is claimed by handle, not by address: a claim says how much shared memory, how many mbarrier slots or how many warps a component needs, never where they are. Second, lifetimes let disjoint uses share storage: a claim may live for the whole kernel, for one phase of a megakernel’s instruction stream, or between two events of a tensor connection. Third, synchronization is described by protocol rather than by code: an mbarrier ring states its arrival counts as functions of the participating roles’ widths, its transaction bytes, and which role initializes it, so that two endpoints can be checked for agreement without inspecting either body.

Table 1. Contract vocabulary: what a component requires from the kernel it joins, and what it promises to the other components. Requires

Promises

Warps for each role Shared memory, barriers, tensor memory, counters, with lifetimes Callables it calls Tensors it imports (demand) Facts it schedules against Caller buffers and communication domains

Roles and their entry functions Channel endpoints and their protocols Callables it exports Tensors it exports (placement) Facts it owns, e.g. tile sizes Host parameters it builds

Inferring these properties from the code instead would require analyzing opaque objects and would still miss intent, such as which tile a buffer carries or when it may be reused. Fourth, facts are typed constants, such as a tile size, that one component owns and others schedule against. A consumer does not compile a fact in: it declares a kernel parameter sourced from the fact, and the linker checks that all declarations agree and writes the owner’s value into the kernel’s parameters, so a changed tile size (Figure 2) reaches every consumer without recompiling it. 3.3

Relocatable Components

A component reads every link-time decision through a handle function: an external device function named after the handle, declared but not defined in the component (Figure 3a, orange). Shared-memory bases, mbarrier slots, named-barrier IDs, warp offsets, ring depths and counter slots are all read this way. The object therefore contains unresolved calls instead of constants, much as positionindependent code reads addresses through a table filled in by the dynamic linker. Unlike a table lookup, a handle function costs nothing at run time: the linker defines it to return a constant that LTO folds away (Section 4.3). We need the object to stay unchanged across compositions, the binding to be chosen per composition, and the result to cost nothing at run time; the obvious alternatives each lose one of these. Template or macro parameters fold into constants but require recompiling the component for every composition, and so its source. Passing placements as kernel arguments or through constant memory keeps the object fixed but turns every offset into a run-time load and blocks constant folding in address arithmetic. Relocation entries patched into the binary, as CPU linkers use, do not exist for most of these resources: named-barrier IDs, warp ranges and register budgets are instruction operands or control flow, not addresses. Handle functions keep the object unchanged and still compile to immediates.

PipeLink: Contract-Guided Linking of GPU Kernel Components

3.4

Authoring

Annotations live in the component’s own source file, next to the code they describe, and are evaluated by the compiler at pack time. The author states only what the compiler cannot see: semantic promises such as which tile an export carries, which protocol a ring follows, which role initializes a barrier, and how long a buffer stays live. Everything the compiler can observe is captured instead of restated. Packing measures the size, alignment and field offsets of parameter and descriptor types, records which device symbols the object defines and needs, reads launch attributes from the compiled object, and fits every enumerated layout map to a closed-form expression, rejecting the package if the fit is not exact. For tensor imports written in TileLang, the frontend captures the consumer’s demand from the layout of its destination fragment, so the author writes no placement facts at all (Figure 3). Optionally, the frontend specializes the consumer to a particular producer by adopting the producer’s published register layout; the result still links with any other producer, but the linker can then choose a direct register copy instead of a staged transfer. Figure 3 compares the two SwiGLU consumers of our GEMM + SwiGLU case study, which import the same accumulator. The CUDA consumer states its contract in 10 device-side annotations spanning 38 lines; the TileLang consumer writes no annotations and 14 lines of frontend calls, which name the export to import and choose the chunk size, specialization and completion policy. The rest is taken from the producer’s contract or captured from the compiled kernel. Both consumers share the same 9 host annotations and host provider; the consumer-specific source is 163 lines in CUDA and 56 in TileLang. Elsewhere in our case studies, a new attention load or communication component takes 13–20 annotation statements and a megakernel operator 11–15; the four owners of the MoE megakernel write 327 statements in total, against 3,482 lines of component and host-provider source.

4

C ONTRACT-G UIDED L INKING

(a) CUDA consumer: device-side contract and its use (abridged) PL STAGE("swiglu.cuda", "1.0.0"); PL TYPE(tile args, a, b, M, N, K, m0, n0); PL ROLE(epilogue, epilogue role, THREADS / 32); PL CALL(tensor produce, ..., pl::produce, ...); PL CALL(tensor release, ..., pl::produce, ...); PL BINDING("chunk", pl::coord{"chunk", 0, 8}, ...); PL SITE(tensor import site, "epilogue", pl::participants{ pl::coord{"t", F, F + 256}, [](int t) { return t; }}, "tile work", pl::guarantees{"enter once", "converged", "full warps", "no early return"}, ...); PL CLAIM SMEM("output smem", OutputTile, ...); PL IMPORT TENSOR(acc, import acc, pl::identity{"accumulator", "tile work", "completed k", "f32", ..., {{"m", 0, 128}, {"n", 0, 256}}}, pl::demand positions{pl::coord{"t", 0, 256}, pl::coord{"p", 0, 16}, pl::bind("chunk"), [](int t, int p, pl::bound c) { return {row(t), 32*c + 16*half row(t) + p, ...}; }}, pl::lends{pl::frame(...), pl::smem("scratch", ...), pl::regs("acc dest", 16, ...), pl::cta barrier(...)}, {pl::rendezvous{"tensor ready", ...}, pl::rendezvous{"import done", ...}}); // kernel code uses the claim and the generated import: auto& tile = *(OutputTile*)(smem + output smem()); import acc(source, chunk, acc, scratch);

(b) TileLang consumer: all it writes for the same contract acc = pl.import tensor(partner, "accumulator", outputs=("m0", "n0"), scratch bytes=0) frag = acc.fragment(specialize=True, chunk=("n", 32)) with acc.produced(completion="final"): staging = acc.staging((128, 128), "bfloat16") for chunk in T.serial(8): with acc.imported(frag, chunk=chunk): ... # SwiGLU on frag, then store staging STAGE = pl.contract stage("swiglu.tilelang", "1.0.0", program, host=...) from the producer’s contract written in both

captured by the frontend uses a claim or the generated import

Figure 3. The same tensor-import contract written by hand in CUDA (a) and through the TileLang frontend (b), abridged. Colors in (a) mark where the TileLang version gets each fact instead. Orange lines use a claim and the generated import; in TileLang the frontend emits the claim’s handle read.

bytes and which warp initializes the barrier. Callables must match their signatures, facts must agree between their owner and their consumers, and topologies such as the cluster shape are intersected across components. Every rejection names one of 44 clauses together with the conflicting facts: a system-scope counter producer paired with a grid-scope consumer, for example, fails with L-PAIR, and a tensormemory claim of 544 columns fails with L-OVERLAP.

A pipeline manifest lists the packages of one composition. Linking it takes four steps (Figure 1): check the contracts, solve the placement, generate the connecting code, and device-link everything with the vendor LTO backend. The output is one kernel image and a record of every decision.

Exports that nothing calls are dropped, as an ordinary linker drops unreferenced sections, so one package can serve several compositions that each call a different subset of its exports.

4.1

A classical linker relocates sections by concatenating them. That does not work on an SM: its resource spaces are small (one thread block gets at most 227 KiB of shared memory, 32 warps and 16 named barriers) and coupled, because the warps a role occupies decide its register group, register

Checking

The linker pairs endpoints by name and checks that each pair agrees on memory space, scope, tile, layout and protocol; for an mbarrier ring, it also checks arrival counts, transaction

4.2

Relocation as Constraint Solving

PipeLink: Contract-Guided Linking of GPU Kernel Components

groups share one register file, and buffers may share storage only when their lifetimes are disjoint. PipeLink instead solves one constraint problem per composition with an SMT solver (de Moura & Bjørner, 2008). The variables are the base of every claimed handle in every domain, the ring depths left to the linker, and the register budget of every register group. The constraints encode capacity, alignment, exact coverage of the warps by roles, and non-overlap. Two claims may overlap only if their phases are mutually exclusive or their lifetimes on one tensor connection are ordered, so an epilogue’s scratch can reuse a mainloop’s operand ring once the accumulator is ready. The objective prefers deeper rings and larger register budgets, then a smaller footprint. The solver is deterministic, so a composition always links to the same placement. A greedy allocator would be simpler, but lifetimes turn placement into packing, where the order of placement decides the footprint. In the MoE megakernel, first-fit packing in lexical or decreasing-size order misses the smallest shared-memory footprint, which the solver finds and which leaves 1,440 of 232,448 bytes free. Our compositions have 4–42 handles each; the solver proves every objective optimal within 1.3% of its work limit. PipeLink also never reschedules a component. Warpspecialization compilers choose pipelining and role assignment for the whole kernel, but they need the source, and the owners of our components have already tuned their schedules for their hardware. Keeping each schedule fixed makes the linked kernel’s performance predictable from its parts and keeps the linker out of the components’ correctness arguments.

role would run at the register budget of the most demanding one. Section 5.1 compares the machine code of linked and single-unit kernels. 4.4

Generated Tensor Transfers

A GEMM mainloop leaves its accumulator in tensor memory or in registers, and an epilogue wants each of its threads to hold particular elements. Writing that transfer by hand requires knowing the producer’s layout, which is the coupling contracts are meant to remove. Instead, the producer’s export publishes the tile’s placement: where each element lives, as a tensor-memory lane and column, a shared-memory address, or a thread and register. The consumer’s import states its demand: which element each of its threads wants (Figure 4). Both are exact maps fitted at pack time. The linker then tries routes from cheapest to most general: a read from tensor memory straight into the demanded registers, a direct register copy when placement and demand coincide, a warp-local shuffle, or staging through scratch shared memory that the consumer lends. It checks each candidate against a model of the hardware primitives for exact delivery of every element, freedom from races and correct ordering, and lowers the first route that passes to CUDA through TileLang (Wang et al., 2026). Transfers move values exactly and never convert them, and their scratch is an ordinary claim.

The vendor backend, not the solver, allocates registers. After the device link, PipeLink reads the linked kernel’s register count and re-solves once if the register groups do not fit; if they still do not fit, the link fails with L-REGS.

Two simpler designs were available. Always staging through shared memory works for every pair, but it costs a round trip that the direct routes avoid (Section 5.5). Fixing one exchange layout by convention would make every producer and consumer agree in advance, which is the coupling contracts remove. At the other extreme, general program synthesis could find cleverer routes, but its results are hard to bound and to check. A short ladder of route shapes, each checked exactly, covers every connection in our case studies.

4.3

4.5

Why Linking Costs Nothing

The linker emits three pieces of code: handle stubs that define every handle function to return its solved constant, call trampolines that connect each call site to the export serving it, and a dispatch kernel that assigns warps to roles and sets each register group’s budget once. All three are compiled to LTO IR and device-linked with the components through nvJitLink (Sundaram et al., 2023). The backend inlines the stubs and folds their constants into the components’ address arithmetic, so a call such as smem base() becomes the immediate 1024. The result is the code a programmer would get by writing the components into one source file with the same placement. Dispatching through function pointers would be simpler to generate, but indirect calls are not inlined, so constants would not fold and every

Communication and Host Interface

Communication stays in component code. A component that moves data between GPUs declares the communication domains it joins, the views it needs of a symmetric allocation the caller lends (peer addresses, a multicast address), and system-scope counters for readiness across GPUs. The linker checks that all owners of an allocation agree on its domain and views, binds counter slots like any other claim, and passes the caller’s addresses to the owners. It does not synthesize collectives: communication protocols are specific to their algorithms (a ring push, a dispatch with routing, a reduce-scatter), and the owners who design them are best placed to implement them. The linker’s job is to make such a component a well-behaved neighbor inside a shared kernel.

PipeLink: Contract-Guided Linking of GPU Kernel Components Placement

Route

color: thread holding it

cheapest that checks

B300 · tensor memory

Demand color: thread wanting it either consumer

Case study

tensor-memory read

H200, RTX 5090 · registers

Table 2. Case studies. Comp. counts distinct component packages; a component is reused if more than one variant links it unchanged.

TileLang consumer

Arch.

Comp. Reused Variants

Megakernel sm103 MoE megakernel sm103 Attention sm90 GEMM + SwiGLU sm90/103/120

10 6 14 9

8 6 8 3

2 24 7 6

register copy

CUDA consumer

1770 MHz on RTX 5090.

shared scratch direct copy

store / load

Figure 4. Generated tensor transfers on the GEMM + SwiGLU case study (schematic). Lines follow single elements from the thread that holds them to the thread that wants them. The CUDA consumer wants elements in other threads, so the direct copy fails the check and the linker stages through shared scratch.

Host-side setup, such as building TMA descriptors, ships as a small compiled provider behind a C interface. When a session opens, each owner’s provider fills its own region of the kernel’s parameters from caller-lent buffers; facts that other owners need flow through the manifest, never through shared headers. An open session launches without host computation and can be captured in a CUDA graph.

5

E VALUATION

We evaluate PipeLink on four case studies (Table 2). Each one asks what a composition makes possible, and each also checks the cost of the composition against the same component sources compiled as a single unit. 5.1

Setup

Implementation. The PipeLink compiler is about 24,000 lines of Python. It packs components with nvcc, solves placement with Z3 (de Moura & Bjørner, 2008), lowers generated transfers through TileLang (Wang et al., 2026), and links with nvJitLink (Sundaram et al., 2023) from CUDA 13.3. A TileLang frontend of about 5,400 lines derives contracts for components written in TileLang. Code and reproduction scripts: https://github.com/PKU -DAIR/PipeLink. Hardware. We use NVIDIA H200 (Hopper, sm90), B300 (Blackwell, sm103) and RTX 5090 (consumer Blackwell, sm120) GPUs. Each machine runs every timed experiment at one locked SM clock, the highest clock it sustains under simultaneous dense BF16 and FP8 matrix multiplication on all GPUs: 1230 MHz on H200, 1155 MHz on B300 and

Methodology. For every linked kernel we build a singleunit control that includes the same component sources, the same generated connecting code and the same placement in one translation unit, compiled with the same floatingpoint mode. Linked and control images run in one process on the same inputs, alternating in balanced blocks with a module reload halfway. Every ratio in this section is a throughput ratio or speedup of the linked kernel, so higher is better. Every linked kernel matches its control (bitwise, or within run-to-run noise for the megakernels) and is free of memcheck, racecheck and synccheck findings. Across the 51 timed configurations, linked kernels run within 0.2% of their single-unit controls. For all seven attention compositions, the mixed-library megakernel, and the GEMM + SwiGLU kernels on H200 and RTX 5090, the normalized machine code of linked and control kernels is identical. The case studies therefore measure what composition enables, not what it costs. 5.2

Cross-Library Megakernel

Composition. We split the Llama-1B (Meta, 2024) decode megakernel of Hazy Research (Spector et al., 2025a), written with ThunderKittens (Spector et al., 2025b), into nine packages: an instruction interpreter with five roles and eight operators. The interpreter runs a precomputed instruction stream for 16 layers in one persistent kernel; operators claim shared-memory pages, semaphores and scratch for the duration of their instruction phase, so different operators reuse the same pool. The original megakernel fails our correctness checks: its split-KV path deadlocks or never stores its partial outputs, and racecheck reports hazards where warps hand shared memory to each other. We repaired five upstream defects in every image, the monolith included. Linked, the all-ThunderKittens kernel reaches 0.9995–1.0002× the throughput of its single-unit control across five configurations, and 1.01–1.10× the throughput of the repaired monolith. Cross-library replacement. We then replace only the attention-partial package with one written in CUTLASS/CuTe (Thakkar et al., 2023): four warps process

PipeLink: Contract-Guided Linking of GPU Kernel Components

Context 16 1024, reduction 4096

TK only 0.713 1.151 3.516

TK + CuTe 0.752 0.939 1.630

causal window

Speedup

softcap

0.95× 1.22× 2.16×

window + softcap Phi-3 block-sparse

context parallel

disjoint key/value tiles with warp-level MMA and merge their normalized partial results within the component. The other eight packages, the instruction stream and the compiled host provider are unchanged; the CuTe package satisfies the same contract, so the composition links without any change to the interpreter. Table 3 shows the result. At a 4096-token context, the mixed-library kernel decodes 2.16× faster than the ThunderKittens kernel; at 1024 tokens with split-KV reduction, 1.22× faster. At a 16-token context it is 5% slower, because one of the four warps has work and the component’s fixed copy and merge costs dominate. The mixed kernel reaches 0.9994–0.9998× the throughput of its own single-unit control and has zero sanitizer findings. 5.3

Multi-Owner MoE Megakernel

Composition. Mixture-of-Kittens (Sul et al., 2026b) is a deterministic expert-parallel MoE training megakernel for Blackwell. We split it among four owners: a framework that owns the task interpreter, the cluster launch control and the launch geometry, and three operator owners for communication (dispatch and combine over NVLink), grouped GEMM, and SwiGLU activation. Each owner builds its own descriptors from caller-lent buffers in a compiled host provider; tile geometry that another owner schedules against is a declared fact, and synchronization between owners runs through counter channels, some with three producers. The framework provides one interpreter package each for the forward, backward and recompute passes; each operator owner ships a single package whose exports serve all three passes and all variants (BF16 and MXFP8, one or four GPUs, clamped or not). These six packages link into all 24 kernels. Performance. On four B300 GPUs at 2048 and 8192 tokens per rank, the 12 timed kernels reach 0.999–1.001× the throughput of their single-unit controls and 0.989–1.014× the throughput of the original monolithic MoK step, measured in the same processes with the same inputs and routing. Owner replacement. An independent activation owner implements the GPT-OSS (OpenAI, 2025) variant of SwiGLU against the same contract. It is built in a source tree that contains none of the other owners’ implementation sources, links into all 24 kernels with the other owners’ pack-

block

MSA sparse CP

core

same compiled package

load

Table 3. Llama-1B decode megakernel on B300 (1155 MHz), ms per token; TK is ThunderKittens.

default, reused

mask

score

epilogue comm.

window softcap window softcap block partial push new

idle warps

Figure 5. Seven attention variants built around one unchanged core. Each row is a linked kernel that links the same compiled core package; the orange cell is the communication component.

ages and providers unchanged, and passes all 48 numerical and bitwise checks per rank. Owner upgrades. All five single-constant upgrades of Section 2.4, which hang or leave outputs unwritten when linked by symbol, relink correctly: PipeLink propagates each new fact to the owners that consume it. Seven incompatible combinations, such as a BF16-only communication package in an MXFP8 kernel, are each rejected with a named diagnostic. 5.4

Attention Variants from One Core

Composition. We derived one attention core for Hopper from FlashAttention-3 (Shah et al., 2024) (BF16, head dimension 128) and froze its package. Seven compositions reuse it unchanged, each adding or replacing at most two components (Figure 5). Sliding windows and softcapping are mask and score callables. Phi-3 (Microsoft, 2024) blocksparse attention and MiniMax sparse attention (MSA) (Lai et al., 2026) replace the load component with one that visits only the selected key blocks. Context-parallel attention (CP) replaces the load component and adds a communication component, which pushes key/value blocks to the other GPUs and releases a readiness counter per block. Because every query tile visits all of its key/value blocks inside one kernel, no output correction is needed. All 17 timed singleGPU configurations reach 0.9991–1.0005× the throughput of their single-unit controls, and CP 0.9997× and 0.9992× on two and four GPUs. Placement. Table 4 compares the variants with public implementations (Shah et al., 2024; Ye et al., 2025) on released-model workloads. The sparse variants gain most, because their load components skip unselected blocks: MSA is 1.6–4.5× faster than vLLM’s (Kwon et al., 2023) dedicated kernel and up to 38× faster than FlexAttention (Dong et al., 2025). FA4 (Zadouri et al., 2026), which masks the selection over query tiles, is 5–6% faster when the selected blocks form a prefix but 9–38× slower when they are scat-

PipeLink: Contract-Guided Linking of GPU Kernel Components Table 4. Attention placement on H200 (1230 MHz): throughput of the linked kernel relative to the best measured alternative. MSA uses one reference for both selection patterns, the alternative fastest over both combined (FA4 is faster on prefixes only). Workload

Reference

Linked / ref.

Qwen3-8B causal, 8K–32K Mistral/Gemma sliding window Gemma2/Grok softcap Llama 4 chunked, 32K–128K Phi-3 block-sparse, 32K–128K MSA, prefix selection MSA, uniform selection

FA3 FA3 FA3 FA3 best expert

0.96–0.99 0.84–0.94 0.92–0.94 1.01–1.02 2.26–2.73 4.47–4.52 1.62–1.63

vLLM Triton

Table 5. Context-parallel causal attention on H200 (1230 MHz), ms. Tokens GPUs PipeLink TE (p2p) Speedup 8K 8K 32K 32K

2 4 2 4

0.604 0.375 7.795 4.133

0.750 0.490 9.446 4.964

cuBLAS + compiled 1.20 / 1.20 1.03 / 1.08 1.03 / 1.01 cuBLAS + Liger 1.34 / 1.34 1.12 / 1.18 1.08 / 1.05 Same GEMM, unfused 1.11 / 1.10 1.02 / 1.07 1.03 / 1.01 Hand-fused native 0.99 / 0.98 0.94 / 0.99 1.01 / 0.99

Table 7. How approaches compose parts of a GPU kernel (WS: warp-specialization; RTs: runtimes; res.: on-chip resources; comm.: cross-GPU communication inside the kernel). • yes, ◦ partly, – no. Binary Cross- Places Checks Kernel parts lib. res. iface. comm. Templates, DSLs WS compilers Megakernel RTs Symbol linking Comm. overlap PipeLink

1.24× 1.31× 1.21× 1.20×

tered.

Context parallelism. For causal prefill over 32K tokens, CP is 1.95× faster on two H200 GPUs and 3.67× faster on four than the single-GPU kernel (97.6% and 91.7% of ideal). Against Transformer Engine’s (NVIDIA, f) context parallelism on the same GPUs, using its faster transport and capturing its forward pass in a CUDA graph, CP is 1.20– 1.31× faster at 8K and 32K tokens (Table 5). Both sides use their own local attention kernels, so the comparison covers complete implementations. 5.5

Table 6. Qwen3-32B GEMM + SwiGLU: speedup of the linked kernel over each baseline at M = 16384; CUDA / TileLang consumer. Baseline B300 H200 5090

Fused GEMM and Activation

Composition. We compose the MLP up-projection of Qwen3-32B (Team, 2025) with its SwiGLU activation: X[M, 5120] times interleaved gate and up weights W [51200, 5120], producing the BF16 activation [M, 25600] for M = 16384 and 4096. The FP32 accumulator stays unrounded through the activation. Splitting a fused kernel into a GEMM producer and an elementwise consumer follows CODA (Guo et al., 2026); our components and shapes differ. Each architecture has its own GEMM producer: a two-SM tensor-core GEMM on B300 that leaves its accumulator in tensor memory, a three-warpgroup WGMMA GEMM on H200 and a multistage mma.sync GEMM on RTX 5090, both leaving it in registers. The two consumers of Section 3.4, one in CUDA and one in TileLang, import the accumulator in 32-column chunks. The producers link without their sources.

– – – • – •

– – – • – •

◦ • ◦ – – •

◦ • ◦ – – •

– – ◦ – • •

Generated transfers. On B300 the linker reads the accumulator from tensor memory into each consumer thread’s registers. On H200 and RTX 5090 it stages each chunk through 16 KiB of shared memory for the CUDA consumer and gives the specialized TileLang consumer a direct register copy (Figure 4). On H200, the direct route makes the TileLang kernel 5% faster than the CUDA one (14.35 versus 15.10 ms at M = 16384). Placement. Table 6 compares the linked kernels with public and native baselines. On B300, the composition is 1.20–1.21× faster than cuBLAS (NVIDIA, a) followed by a torch.compiled (Ansel et al., 2024) SwiGLU and 1.34–1.37× faster than cuBLAS followed by Liger’s (Hsu et al., 2024) SwiGLU. Against fusing the same components by hand into one native kernel, the linked kernels reach 0.94–1.01× its throughput. All 12 configurations reach 0.9987–1.0015× the throughput of their single-unit controls.

6

R ELATED W ORK

We group prior work by where the information needed to compose kernel parts lives (Table 7). In one source program. Template libraries such as CUTLASS (Thakkar et al., 2023) and attention libraries that accept mask and score functions or generate variants from templates (Dong et al., 2025; Ye et al., 2025; Chen et al., 2025; You et al., 2026; Chen et al., 2026a) compose at extension points that their authors designed. Domain-specific

PipeLink: Contract-Guided Linking of GPU Kernel Components

languages such as TileLang (Wang et al., 2026), CuTe DSL (NVIDIA, d), ThunderKittens (Spector et al., 2025b) and Triton (Tillet et al., 2019), and fusion libraries such as FKL (Amoros et al., 2025) and CODA (Guo et al., 2026), compose within one program; sparse-attention kernels specialize execution to the selection pattern in the same way (Zhao et al., 2026; Li et al., 2026b; Liu et al., 2026). PipeLink instead links components that different authors compiled separately, and can link components written with these tools. In compiler IR. Warp-specialization compilers partition a kernel into roles and synchronize them: Tawa (Chen et al., 2026b) uses asynchronous references as channels, Twill (Soi et al., 2026) solves pipelining and warp specialization jointly, Cypress (Yadav et al., 2025) derives data movement from a mapping, and TLX (Guan et al., 2026) and Gluon (Triton contributors) expose roles and memory to the programmer. Layout systems (Zhang et al., 2026; Zhou et al., 2026; Hou et al., 2026) represent the layouts that connect operations. These compilers own and may reschedule the whole kernel; PipeLink keeps each component’s compiled schedule and works over declared interfaces of opaque objects. In runtime task graphs. Megakernel systems schedule operators or tiles across SMs in one persistent kernel: MPK (Cheng et al., 2026) and Event Tensor (Jin et al., 2026) compile models into in-kernel task graphs, and other work tunes, searches or generates megakernel schedules, also with coding agents (Spector et al., 2025a; He et al., 2026a;b; Jaber & Jaber, 2026; Dong et al., 2026; Chowdhary et al., 2026; Li et al., 2026a; Yao et al., 2026). These systems decide which task runs where; PipeLink decides how separately compiled components share one SM, so their task bodies could be linked components, as in our Llama-1B megakernel. By symbol. Relocatable device code and LTO (Scudiero & Murphy, 2014), exposed at run time by nvJitLink (Sundaram et al., 2023), link separately compiled device functions at no runtime cost; cuFFT (NVIDIA, c) and cuFFTDx (NVIDIA, b) use them to fuse callbacks and library code, and GoPTX (Wu et al., 2025) fuses kernels at the PTX level. PipeLink adds what concurrently running components share and a linker that checks and places it. Communication overlap. Many systems overlap communication with computation through streams, signals or fused kernels (Chang et al., 2024; Zheng et al., 2025b;a; Awad et al., 2025; Sul et al., 2026a; Qiang et al., 2026; Mao et al., 2026; Ma et al., 2026; Hong et al., 2026; Gond et al., 2026; Zhang et al., 2025; NVIDIA, e), and contextparallel frameworks correct outputs outside the attention

kernel (NVIDIA, f; Fang & Zhao, 2024; Gu et al., 2024; Yang et al., 2025). PipeLink contributes no communication algorithm; it lets one owner’s communication component join another owner’s compute components in one kernel.

Synchronization verification. WEFT (Sharma et al., 2015), GPUVerify (Betts et al., 2012), Descend (Köpcke et al., 2024) and Exo-GPU (Akeley et al., 2026) verify synchronization or schedule safety of whole programs. PipeLink checks declared protocols between components and leaves their code to sanitizers and numerical tests.

7

L IMITATIONS AND C ONCLUSION

Limitations. PipeLink checks contracts against each other, not against the code behind them. A component that breaks its own contract, for example by touching shared memory outside its claim, can still produce a wrong kernel; numerical checks and sanitizers therefore remain part of every composition, as they are for any hand-fused kernel. The code PipeLink generates only moves data, and only within one thread block. Arithmetic, precision conversion and communication between thread blocks or GPUs stay in component code, where the linker checks and binds the counters and memory they use but does not write them. A tensor export carries one generation, so a producer that hands its accumulator to another component in this way cannot be persistent; multi-slot rings between components are written in component code. PipeLink does not make components faster. It preserves each owner’s schedule, so a composition performs as well as the components it links; the gap of our dense attention variants to FlashAttention-3 comes from our core. Existing kernels must also be ported to declared interfaces; the fourowner MoE megakernel shows what such a port involves. Finally, packages are compiled for one architecture and one toolkit, because LTO IR is not portable across CUDA major versions.

Conclusion. The components of fused GPU kernels cooperate at fine grain on one SM, yet today they are absoluteaddressed and can only be combined by forking their source. PipeLink gives them what object files gave whole programs: a contract stating what crosses the component boundary, and a linker that checks the contracts and relocates every component into the kernel it joins. Because the linker’s decisions fold into constants, composition costs nothing, and components from different libraries, owners and GPUs combine into kernels that were previously built by hand or not at all.

PipeLink: Contract-Guided Linking of GPU Kernel Components

R EFERENCES Akeley, D. Z., Ikarashi, Y., and Ragan-Kelley, J. Exo-GPU: Safe, Imperative, User-schedulable Programming for Tensor Cores. https://arxiv.org/abs/2609.163 89, 2026. Accessed 2026-10-20. Amoros, O., Andaluz, A., Nunez, J., and Pena, A. J. The fused kernel library: A c++ API to develop highlyefficient GPU libraries. CoRR, abs/2508.07071, 2025. doi: 10.48550/ARXIV.2508.07071. URL https: //doi.org/10.48550/arXiv.2508.07071. Ansel, J., Yang, E. Z., He, H., Gimelshein, N., Jain, A., Voznesensky, M., Bao, B., Bell, P., Berard, D., Burovski, E., Chauhan, G., Chourdia, A., Constable, W., Desmaison, A., DeVito, Z., Ellison, E., Feng, W., Gong, J., Gschwind, M., Hirsh, B., Huang, S., Kalambarkar, K., Kirsch, L., Lazos, M., Lezcano, M., Liang, Y., Liang, J., Lu, Y., Luk, C., Maher, B., Pan, Y., Puhrsch, C., Reso, M., Saroufim, M., Siraichi, M. Y., Suk, H., Zhang, S., Suo, M., Tillet, P., Zhao, X., Wang, E., Zhou, K., Zou, R., Wang, X., Mathews, A., Wen, W., Chanan, G., Wu, P., and Chintala, S. PyTorch 2: Faster machine learning through dynamic python bytecode transformation and graph compilation. In Gupta, R., Abu-Ghazaleh, N. B., Musuvathi, M., and Tsafrir, D. (eds.), Proceedings of the 29th ACM International Conference on Architectural Support for Programming Languages and Operating Systems, Volume 2, ASPLOS 2024, La Jolla, CA, USA, 27 April 2024- 1 May 2024, pp. 929–947. ACM, 2024. doi: 10.1145/3620665.3640366. URL https: //doi.org/10.1145/3620665.3640366. Awad, M., Osama, M., and Potter, B. Iris: First-class multi-GPU programming experience in triton. CoRR, abs/2511.12500, 2025. doi: 10.48550/ARXIV.2511.12 500. URL https://doi.org/10.48550/arXiv .2511.12500. Betts, A., Chong, N., Donaldson, A. F., Qadeer, S., and Thomson, P. GPUVerify: a verifier for GPU kernels. In Leavens, G. T. and Dwyer, M. B. (eds.), Proceedings of the 27th Annual ACM SIGPLAN Conference on Object-Oriented Programming, Systems, Languages, and Applications, OOPSLA 2012, part of SPLASH 2012, Tucson, AZ, USA, October 21-25, 2012, pp. 113–132. ACM, 2012. doi: 10.1145/2384616.2384625. URL https: //doi.org/10.1145/2384616.2384625. Chang, L., Bao, W., Hou, Q., Jiang, C., Zheng, N., Zhong, Y., Zhang, X., Song, Z., Jiang, Z., Lin, H., Jin, X., and Liu, X. FLUX: fast software-based communication overlap on GPUs through kernel fusion. CoRR, abs/2406.06858, 2024. doi: 10.48550/ARXIV.2406.06858. URL https: //doi.org/10.48550/arXiv.2406.06858.

Chen, F., Cheng, Y., Wang, L., Xia, Y., Miao, Z., Ma, L., Yang, F., Xue, J., Yang, Z., Yang, M., and Chen, H. AttentionEngine: A versatile framework for efficient attention mechanisms on diverse hardware platforms. CoRR, abs/2502.15349, 2025. doi: 10.48550/ARXIV.2502.15 349. URL https://doi.org/10.48550/arXiv .2502.15349. Chen, F., Cheng, Y., Wang, L., Xia, Y., Miao, Z., Ma, L., Yang, F., Xue, J., Yang, Z., Yang, M., Wei, X., and Chen, H. MetaAttention: A unified and performant attention framework across hardware backends. In Hosking, T., Musuvathi, M., and Taura, K. (eds.), Proceedings of the 31st ACM SIGPLAN Annual Symposium on Principles and Practice of Parallel Programming, PPoPP 2026, Sydney, NSW, Australia, 31 January 2026 - 4 February 2026, pp. 635–647. ACM, 2026a. doi: 10.1145/3774934.3786444. URL https: //doi.org/10.1145/3774934.3786444. Chen, H., Fan, B., Collins, A., Hagedorn, B., Gaburov, E., Masuda, M., Brookhart, M., Sullivan, C., Knight, J., Zhang, Z., and Grover, V. Tawa: Automatic warp specialization for modern GPUs with asynchronous references. In Blackburn, S. M., Cohen, A., and Jones, T. M. (eds.), IEEE/ACM International Symposium on Code Generation and Optimization, CGO 2026, Sydney, Australia, January 31 - Feb. 4, 2026, pp. 255–267. IEEE, 2026b. doi: 10.1109/CGO68049.2026.11394849. URL https://doi.org/10.1109/CGO68049.2026. 11394849. Cheng, X., Zhang, Z., Zhou, Y., Ji, J., Jiang, J., Zhao, Z., Xiao, Z., Ye, Z., Huang, Y., Lai, R., Jin, H., Hou, B., Wu, M., Dong, Y., Yip, A., Ye, Z., Wang, S., Yang, W., Miao, X., Chen, T., and Jia, Z. MPK: A Compiler and Runtime for Mega-Kernelizing Tensor Programs. In 20th USENIX Symposium on Operating Systems Design and Implementation (OSDI 26), 2026. URL https://www. usenix.org/conference/osdi26/present ation/cheng. Chowdhary, S., Swann, R., Siddens, S., Osama, M., Neuendorffer, S., Dutu, A., Sangaiah, K., Bhuyan, S., Bayliss, S., and Dasika, G. Fleet: Hierarchical task-based abstraction for megakernels on multi-die GPUs. CoRR, abs/2604.15379, 2026. doi: 10.48550/ARXIV.2604.15 379. URL https://doi.org/10.48550/arXiv .2604.15379. de Moura, L. M. and Bjørner, N. S. Z3: an efficient SMT solver. In Ramakrishnan, C. R. and Rehof, J. (eds.), Tools and Algorithms for the Construction and Analysis of Systems, 14th International Conference, TACAS 2008, Held as Part of the Joint European Conferences on Theory and Practice of Software, ETAPS 2008, Budapest,

PipeLink: Contract-Guided Linking of GPU Kernel Components

Hungary, March 29-April 6, 2008. Proceedings, volume 4963 of Lecture Notes in Computer Science, pp. 337–340. Springer, 2008. doi: 10.1007/978-3-540-78800-3\ 24. URL https://doi.org/10.1007/978-3-540 -78800-3_24. Dong, J., Feng, B., Guessous, D., Liang, Y., and He, H. FlexAttention: A programming model for generating fused attention variants. In Zaharia, M., Joshi, G., and Lin, Y. C. (eds.), Proceedings of the Eighth Conference on Machine Learning and Systems, MLSys 2025, Santa Clara, CA, USA, May 12-15, 2025. OpenReview.net/mlsys.org, 2025. URL https://openreview.net/forum?id= 2QMYV4bA0R. Dong, W., Hu, M., Yu, G., Fu, Q., Xu, P., Xu, H., Xing, Y., Jiao, X., Li, S., and Liu, L. Ada-MK: Adaptive MegaKernel optimization via automated DAG-based search for LLM inference. CoRR, abs/2605.11581, 2026. doi: 10.48550/ARXIV.2605.11581. URL https: //doi.org/10.48550/arXiv.2605.11581. Fang, J. and Zhao, S. USP: A unified sequence parallelism approach for long context generative AI. CoRR, abs/2405.07719, 2024. doi: 10.48550/ARXIV.2405.0771 9. URL https://doi.org/10.48550/arXiv.2 405.07719. Gond, R., Kwatra, N., and Ramjee, R. TokenWeave: Efficient Compute-Communication Overlap for Distributed LLM Inference. In Proceedings of Machine Learning and Systems (MLSys), 2026. URL https://procee dings.mlsys.org/paper_files/paper/20 26/hash/73ba81c7b25134a559c8a9c39ec1 a4c3-Abstract-Conference.html. Gu, D., Sun, P., Hu, Q., Huang, T., Chen, X., Xiong, Y., Wang, G., Chen, Q., Zhao, S., Fang, J., Wen, Y., Zhang, T., Jin, X., and Liu, X. LoongTrain: Efficient training of long-sequence LLMs with head-context parallelism. CoRR, abs/2406.18485, 2024. doi: 10.48550/ARXIV.2 406.18485. URL https://doi.org/10.48550/a rXiv.2406.18485. Guan, Y., Yu, H., Chen, P., Shi, D., Manivannan, K., Riasanovsky, N. J., Ren, M., Wang, L., Nay, S., Kanuparthy, P., Pan, Z., Hu, Z., and Ding, Y. TLX: hardwarenative, evolvable MIMW GPU compiler for large-scale production environments. CoRR, abs/2605.10905, 2026. doi: 10.48550/ARXIV.2605.10905. URL https: //doi.org/10.48550/arXiv.2605.10905. Guo, H., Zhang, J., Menon, A., Guessous, D., Thakkar, V., Kim, Y., and Dao, T. CODA: Rewriting Transformer Blocks as GEMM-Epilogue Programs. In Advances in Neural Information Processing Systems (NeurIPS), 2026.

URL https://neurips.cc/virtual/2026/ loc/atlanta/poster/153802. He, J., Sun, G., Li, T., Du, D., Xia, Y., and Chen, H. Taming Dynamism on GPUs: Cross-SM Kernel Fusion via SM Cooperation and Just-in-Time Reduction. https:// doi.org/10.1145/3830418.3843857, 2026a. Accessed 2026-10-20. He, Z., Sampson, A., Zhang, Y., and Guo, Z. VDCores: Resource decoupled programming and execution for asynchronous GPU. CoRR, abs/2605.03190, 2026b. doi: 10.48550/ARXIV.2605.03190. URL https: //doi.org/10.48550/arXiv.2605.03190. Hong, K., Li, X., Liu, M., Mao, Q., Wu, T., Huang, Z., Chen, L., Wang, Z., Zhang, Y., Zhu, Z., Dai, G., and Wang, Y. Efficient and adaptable overlapping for computation and communication via signaling and reordering. In Barbalace, A., Mai, L., Geambasu, R., and Pietzuch, P. R. (eds.), Proceedings of the 21st European Conference on Computer Systems, EuroSys 2026, McEwan Hall/The University of Edinburgh, Edinburgh, Scotland, UK, April 27-30, 2026, pp. 1894–1911. ACM, 2026. doi: 10.1145/3767295.3769370. URL https: //doi.org/10.1145/3767295.3769370. Hou, B., Jin, H., Wang, G., Chen, J., Cai, Y., Yang, L., Ye, Z., Ding, Y., Lai, R., and Chen, T. Axe: A simple unified layout abstraction for machine learning compilers. CoRR, abs/2601.19092, 2026. doi: 10.48550/ARXIV.2601.19 092. URL https://doi.org/10.48550/arXiv .2601.19092. Hsu, P., Dai, Y., Kothapalli, V., Song, Q., Tang, S., Zhu, S., Shimizu, S., Sahni, S., Ning, H., and Chen, Y. Liger kernel: Efficient triton kernels for LLM training. CoRR, abs/2410.10989, 2024. doi: 10.48550/ARXIV.2410.10 989. URL https://doi.org/10.48550/arXiv .2410.10989. Hu, B., V, Y. S. V., Agarwal, S., and Akella, A. CUCo: An agentic framework for compute and communication co-design. CoRR, abs/2603.02376, 2026. doi: 10.48550 /ARXIV.2603.02376. URL https://doi.org/10 .48550/arXiv.2603.02376. Jaber, J. and Jaber, O. AutoMegaKernel: A staticallychecked agent harness for self-retargeting megakernel synthesis. CoRR, abs/2606.09682, 2026. doi: 10.48550 /ARXIV.2606.09682. URL https://doi.org/10 .48550/arXiv.2606.09682. Jin, H., Hou, B., Wang, G., Lai, R., Chen, J., Ye, Z., Cai, Y., Dong, Y., Cheng, X., Zhang, Z., Zhao, Y., Huang, Y., Yang, L., Jiang, J., Oliaro, G., Ji, J., Miao, X., Grover, V.,

PipeLink: Contract-Guided Linking of GPU Kernel Components

Mowry, T. C., Jia, Z., and Chen, T. Event Tensor: A Unified Abstraction for Compiling Dynamic Megakernel. In Proceedings of Machine Learning and Systems (MLSys), 2026. URL https://proceedings.mlsys.or g/paper_files/paper/2026/hash/53d3f4 5797970d323bd8a0d379c525aa-Abstract-C onference.html. Köpcke, B., Gorlatch, S., and Steuwer, M. Descend: A safe GPU systems programming language. Proc. ACM Program. Lang., 8(PLDI):841–864, 2024. doi: 10.1145/ 3656411. URL https://doi.org/10.1145/36 56411. Kwon, W., Li, Z., Zhuang, S., Sheng, Y., Zheng, L., Yu, C. H., Gonzalez, J., Zhang, H., and Stoica, I. Efficient memory management for large language model serving with PagedAttention. In Flinn, J., Seltzer, M. I., Druschel, P., Kaufmann, A., and Mace, J. (eds.), Proceedings of the 29th Symposium on Operating Systems Principles, SOSP 2023, Koblenz, Germany, October 23-26, 2023, pp. 611– 626. ACM, 2023. doi: 10.1145/3600006.3613165. URL https://doi.org/10.1145/3600006.3613 165. Lai, X., Xu, W., Yang, Y., Chen, Q., Xu, Y., Zeng, L., Li, X., Sun, H., Zhu, H., Zhang, V., Hu, J., Li, J., Gao, R., Li, Z., Zhu, S., Zhou, J., and Zhao, P. MiniMax sparse attention. CoRR, abs/2606.13392, 2026. doi: 10.48550 /ARXIV.2606.13392. URL https://doi.org/10 .48550/arXiv.2606.13392. Li, L., Zhu, Z., Deng, X., Chen, Y., He, Q., Li, Y., Zhao, R., Han, X., and Liu, Z. ForgeMegakernel: A General Framework for Efficient Auto-Regressive Model Decode Megakernels. https://arxiv.org/abs/2609.1 2379, 2026a. Accessed 2026-10-20. Li, Z., Li, Z., Yuan, Z., Wang, L., and Wang, H. PageWeaver: KV-Guided Query Unions for Sparse Attention. https: //arxiv.org/abs/2610.11201, 2026b. Accessed 2026-10-20. Liu, S., Yu, X., Li, W., and Jiang, W. Decoupling Logical Masks from GPU Execution for Dynamic Block-Sparse Attention. https://arxiv.org/abs/2609.258 69, 2026. Accessed 2026-10-20. Ma, K., Lv, Q., Ge, J., Dai, B., and Ruan, K. Entwine: Coordinating tiled computation and fine-grained communication across GPUs. https://arxiv.org/abs/ 2609.11562, 2026. Accessed 2026-10-20. Mao, Z., Zhang, Y., Chew, S. W., Ma, S., Raiciu, C., Zhou, Y., Shenker, S., and Stoica, I. mKernel: Fast Multi-GPU, Multi-Node fused kernels. https://arxiv.org/ abs/2609.13585, 2026. Accessed 2026-10-20.

Meta. Llama 3.2 Model Card. https://github.c om/meta-llama/llama-models/blob/0c2 12220b07b005201d268e3de1685bf3932841 2/models/llama3_2/MODEL_CARD.md, 2024. Accessed 2026-10-20. Microsoft. Phi-3 technical report: A highly capable language model locally on your phone. CoRR, abs/2404.14219, 2024. doi: 10.48550/ARXIV.2404. 14219. URL https://doi.org/10.48550/arX iv.2404.14219. NVIDIA. cuBLAS: The CUDA Basic Linear Algebra Subroutine Library. https://docs.nvidia.com/cu da/cublas/index.html, a. Accessed 2026-10-20. NVIDIA. Using cuFFTDx with cuFFT LTO. https: //docs.nvidia.com/cuda/cufftdx/cufft _lto/index.html, b. Accessed 2026-10-20. NVIDIA. NVIDIA cuFFT LTO EA Preview Documentation. https://docs.nvidia.com/cuda/arc hive/12.4.1/cufft/ltoea/index.html, c. Accessed 2026-10-20. NVIDIA. CuTe DSL: Introduction. https://docs.n vidia.com/cutlass/latest/media/docs/ pythonDSL/cute_dsl_general/dsl_intro duction.html, d. Accessed 2026-10-20. NVIDIA. NVIDIA OpenSHMEM Library (NVSHMEM) Documentation. https://docs.nvidia.com/nv shmem/api/index.html, e. Accessed 2026-10-20. NVIDIA. Transformer Engine. https://github.c om/NVIDIA/TransformerEngine, f. Accessed 2026-10-20. OpenAI. gpt-oss-120b & gpt-oss-20b model card. CoRR, abs/2508.10925, 2025. doi: 10.48550/ARXIV.2508.10 925. URL https://doi.org/10.48550/arXiv .2508.10925. Qiang, X., Guan, Y., Hu, Z., Zhou, K., Ding, Y., and Aziz, A. Syncopate: Efficient Multi-GPU AI Kernels via Automatic Chunk-Centric Compute-Communication Overlap. In 20th USENIX Symposium on Operating Systems Design and Implementation (OSDI 26), 2026. URL https://www.usenix.org/conference/os di26/presentation/qiang. Scudiero, T. and Murphy, M. Separate Compilation and Linking of CUDA C++ Device Code. https://deve loper.nvidia.com/blog/separate-compi lation-linking-cuda-device-code/, 2014. Accessed 2026-10-20.

PipeLink: Contract-Guided Linking of GPU Kernel Components

Shah, J., Bikshandi, G., Zhang, Y., Thakkar, V., Ramani, P., and Dao, T. FlashAttention-3: Fast and accurate attention with asynchrony and low-precision. In Globersons, A., Mackey, L., Belgrave, D., Fan, A., Paquet, U., Tomczak, J. M., and Zhang, C. (eds.), Advances in Neural Information Processing Systems 37: Annual Conference on Neural Information Processing Systems 2024, NeurIPS 2024, Vancouver, BC, Canada, December 10 - 15, 2024, 2024. URL http://papers.nips.cc/paper_f iles/paper/2024/hash/7ede97c3e082c6d f10a8d6103a2eebd2-Abstract-Conference. html. Sharma, R., Bauer, M., and Aiken, A. Verification of producer-consumer synchronization in GPU programs. In Grove, D. and Blackburn, S. M. (eds.), Proceedings of the 36th ACM SIGPLAN Conference on Programming Language Design and Implementation, Portland, OR, USA, June 15-17, 2015, pp. 88–98. ACM, 2015. doi: 10.1145/2737924.2737962. URL https: //doi.org/10.1145/2737924.2737962. Soi, R., Yadav, R., Kjolstad, F., Aiken, A., Dehnavi, M. M., Garland, M., and Bauer, M. Optimal Software Pipelining and Warp Specialization for Tensor Core GPUs. In 20th USENIX Symposium on Operating Systems Design and Implementation (OSDI 26), 2026. URL https://www. usenix.org/conference/osdi26/present ation/soi. Spector, B., Juravsky, J., Sul, S., Dugan, O., Lim, D., Fu, D., Arora, S., and Ré, C. Look Ma, No Bubbles! Designing a Low-Latency Megakernel for Llama-1B. https:// hazyresearch.stanford.edu/blog/2025-0 5-27-no-bubbles, 2025a. Accessed 2026-10-20. Spector, B. F., Arora, S., Singhal, A., Parthasarathy, A., Fu, D. Y., and Ré, C. ThunderKittens: Simple, fast, and adorable kernels. In The Thirteenth International Conference on Learning Representations, ICLR 2025, Singapore, April 24-28, 2025. OpenReview.net, 2025b. URL https://openreview.net/forum?id= 0fJfVOSUra. Sul, S. H., Arora, S., Spector, B. F., and Ré, C. ParallelKittens: Systematic and Practical Simplification of MultiGPU AI Kernels. In Proceedings of Machine Learning and Systems (MLSys), 2026a. URL https://procee dings.mlsys.org/paper_files/paper/20 26/hash/ff997469ac66cf893c4183efeb22 212a-Abstract-Conference.html. Sul, S. H., Brown, N., Wildermuth, H., Lin, W., Cassano, F., and Ré, C. Mixture-of-Kittens: MoE Megakernel for NVL72s. https://arxiv.org/abs/2609.360 70, 2026b. Accessed 2026-10-20.

Sundaram, A., Murphy, M., Avila, M. F., Ligowski, L., and Oh, F. CUDA 12.0 Compiler Support for Runtime LTO Using nvJitLink Library. https://developer.nv idia.com/blog/cuda-12-0-compiler-sup port-for-runtime-lto-using-nvjitlink -library/, 2023. Accessed 2026-10-20. Team, Q. Qwen3 technical report. CoRR, abs/2505.09388, 2025. doi: 10.48550/ARXIV.2505.09388. URL https: //doi.org/10.48550/arXiv.2505.09388. Thakkar, V., Ramani, P., Cecka, C., Shivam, A., Lu, H., Yan, E., Kosaian, J., Hoemmen, M., Wu, H., Kerr, A., Nicely, M., Merrill, D., Blasig, D., Atluri, A., Qiao, F., Majcher, P., Springer, P., Hohnerbach, M., Wang, J., and Gupta, M. CUTLASS. https://github.com/NVIDIA/cu tlass, 2023. Accessed 2026-10-20. Tillet, P., Kung, H., and Cox, D. D. Triton: an intermediate language and compiler for tiled neural network computations. In Mattson, T., Muzahid, A., and Solar-Lezama, A. (eds.), Proceedings of the 3rd ACM SIGPLAN International Workshop on Machine Learning and Programming Languages, MAPL@PLDI 2019, Phoenix, AZ, USA, June 22, 2019, pp. 10–19. ACM, 2019. doi: 10.1145/3315508.3329973. URL https: //doi.org/10.1145/3315508.3329973. Triton contributors. Gluon Overview. https://trit on-lang.org/main/gluon/index.html. Accessed 2026-10-20. Wang, L., Cheng, Y., Shi, Y., Mo, Z., Tang, Z., Xie, W., Wu, T., Ma, L., Xia, Y., Xue, J., Yang, F., and Yang, Z. TileLang: Bridge programmability and performance in modern neural kernels. In Vondrick, C., Hariharan, B., Raffel, C., Pinto, L., Yang, D., and Faust, A. (eds.), The Fourteenth International Conference on Learning Representations, ICLR 2026, Rio de Janeiro, Brazil, April 23-27, 2026. proceedings.iclr.cc, 2026. URL https: //proceedings.iclr.cc/paper_files/pa per/2026/hash/76fb92288bf90360c527ef b0d1c2aba6-Abstract-Conference.html. Wu, K., Lin, Z., Xi, M., Zheng, Z., Pan, W., Zhang, X., and Lu, Y. GoPTX: Fine-grained GPU kernel fusion by PTX-level instruction flow weaving. In 62nd ACM/IEEE Design Automation Conference, DAC 2025, San Francisco, CA, USA, June 22-25, 2025, pp. 1–7. IEEE, 2025. doi: 10.1109/DAC63849.2025.11132627. URL https://doi.org/10.1109/DAC63849.2025. 11132627. Yadav, R., Garland, M., Aiken, A., and Bauer, M. Taskbased tensor computations on modern GPUs. Proc. ACM Program. Lang., 9(PLDI):396–420, 2025. doi: 10.1145/

PipeLink: Contract-Guided Linking of GPU Kernel Components

3729262. URL https://doi.org/10.1145/37 29262. Yang, A., Yang, J., Ibrahim, A., Xie, X., Tang, B., Sizov, G., Park, J., and Huang, J. Context parallelism for scalable million-token inference. In Zaharia, M., Joshi, G., and Lin, Y. C. (eds.), Proceedings of the Eighth Conference on Machine Learning and Systems, MLSys 2025, Santa Clara, CA, USA, May 12-15, 2025. OpenReview.net/mlsys.org, 2025. URL https://openrevi ew.net/forum?id=Vmf09yVJhT. Yao, J., Hari, S. K. S., Balaji, V., Damani, S., Jang, I., Viswanath, P., and Kozyrakis, C. MegaFlux: SkewResilient MoE Megakernels via Pipelined Expert Replication. https://arxiv.org/abs/2610.00671, 2026. Accessed 2026-10-20. Ye, Z., Chen, L., Lai, R., Lin, W., Zhang, Y., Wang, S., Chen, T., Kasikci, B., Grover, V., Krishnamurthy, A., and Ceze, L. FlashInfer: Efficient and customizable attention engine for LLM inference serving. In Zaharia, M., Joshi, G., and Lin, Y. C. (eds.), Proceedings of the Eighth Conference on Machine Learning and Systems, MLSys 2025, Santa Clara, CA, USA, May 12-15, 2025. OpenReview.net/mlsys.org, 2025. URL https://op enreview.net/forum?id=RXPofAsL8F. You, B., Wang, I., Mustafaoglu, Z. S., Jangda, A., Moreira, A., Dathathri, R., Mahajan, D., and Pingali, K. Flashlight: PyTorch Compiler Extensions to Accelerate Attention Variants. In Proceedings of Machine Learning and Systems (MLSys), 2026. URL https://proceedings. mlsys.org/paper_files/paper/2026/has h/bc52716d13d2d72ea0f335667d86c0f8-A bstract-Conference.html.

Zhang, X., Ding, Y., Sun, B., Hu, Y., Shpeisman, T., and Pekhimenko, G. Hexcute: A compiler framework for automating layout synthesis in GPU programs. In Blackburn, S. M., Cohen, A., and Jones, T. M. (eds.), IEEE/ACM International Symposium on Code Generation and Optimization, CGO 2026, Sydney, Australia, January 31 - Feb. 4, 2026, pp. 630–643. IEEE, 2026. doi: 10.1109/CGO68049.2026.11395194. URL https://doi.org/10.1109/CGO68049.2026. 11395194. Zhao, Z., Zhou, Q., Song, M., Chen, Z., Ye, C., Hou, F., Gong, Z., Li, J., Si, H., and Long, G. QUILT: Rethinking Sparse-Attention Prefill through Shared Query Execution. https://arxiv.org/abs/2610.11134, 2026. Accessed 2026-10-20. Zheng, S., Bao, W., Hou, Q., Zheng, X., Fang, J., Huang, C., Li, T., Duanmu, H., Chen, R., Xu, R., Guo, Y., Zheng, N., Jiang, Z., Di, X., Wang, D., Ye, J., Lin, H., Chang, L., Lu, L., Liang, Y., Zhai, J., and Liu, X. Triton-distributed: Programming overlapping kernels on distributed AI systems with the triton compiler. CoRR, abs/2504.19442, 2025a. doi: 10.48550/ARXIV.2504.19442. URL https: //doi.org/10.48550/arXiv.2504.19442. Zheng, S., Fang, J., Zheng, X., Hou, Q., Bao, W., Zheng, N., Jiang, Z., Wang, D., Ye, J., Lin, H., Chang, L., and Liu, X. TileLink: Generating efficient compute-communication overlapping kernels using tile-centric primitives. In Zaharia, M., Joshi, G., and Lin, Y. C. (eds.), Proceedings of the Eighth Conference on Machine Learning and Systems, MLSys 2025, Santa Clara, CA, USA, May 12-15, 2025. OpenReview.net/mlsys.org, 2025b. URL https: //openreview.net/forum?id=ccjvBkTRRe.

Zadouri, T., Hoehnerbach, M., Shah, J., Thakkar, V., and Dao, T. FlashAttention-4: Algorithm and Kernel Pipelining Co-Design for Asymmetric Hardware Scaling. In Proceedings of Machine Learning and Systems (MLSys), 2026. URL https://proceedings.mlsys.or g/paper_files/paper/2026/hash/ae8b0b 5838ba510daff1198474e7b984-Abstract-C onference.html.

Zheng, S., Zheng, X., Chang, L., and Zhai, J. UniEP: Unified expert-parallel MegaKernel MoE for LLM training. In Kuppannagari, S., Koyutürk, M., Goldman, A., and Nikolopoulos, D. S. (eds.), Proceedings of the 35th International Symposium on High-Performance Parallel and Distributed Computing, HPDC 2026, Cleveland, OH, USA, July 13-16, 2026, pp. 387–401. ACM, 2026. doi: 10.1145/3806645.3807818. URL https: //doi.org/10.1145/3806645.3807818.

Zhang, S., Zheng, N., Lin, H., Jiang, Z., Bao, W., Jiang, C., Hou, Q., Cui, W., Zheng, S., Chang, L., Chen, Q., and Liu, X. COMET: fine-grained computationcommunication overlapping for mixture-of-experts. In Zaharia, M., Joshi, G., and Lin, Y. C. (eds.), Proceedings of the Eighth Conference on Machine Learning and Systems, MLSys 2025, Santa Clara, CA, USA, May 12-15, 2025. OpenReview.net/mlsys.org, 2025. URL https: //openreview.net/forum?id=fGgQS5VW09.

Zhou, K., Casado, M. L., Goucher, A. P., Rakhmati, A., Niu, J., Lebar, J., Szczerbuk, P., Bell, P., Tillet, P., Raoux, T., and Moudallal, Z. Linear layouts: Robust code generation of efficient tensor computation using f 2. In Lee, B. C., Xu, H., Silberstein, M., and Li, B. (eds.), Proceedings of the 31st ACM International Conference on Architectural Support for Programming Languages and Operating Systems, Volume 1, ASPLOS 2026, Pittsburgh, PA, USA, March 22-26, 2026, pp. 132–146. ACM,

PipeLink: Contract-Guided Linking of GPU Kernel Components

2026. doi: 10.1145/3760250.3762221. URL https: //doi.org/10.1145/3760250.3762221. zhuzilin. Ring Flash Attention. https://github.com /zhuzilin/ring-flash-attention. Accessed 2026-10-20.

Related documents

• PipeLink: Contract-Guided Linking of GPU Kernel Components#1468804
• KernelZero: Co-Evolving Proposer and Coder for Continuously Improved GPU Kernel Generation#1117786
• mKernel: Fast Multi-GPU, Multi-Node Fused Kernels#1071362
• mKernel: Fast Multi-GPU, Multi-Node Fused Kernels#919318
• PipeMax: Enhancing Offline LLM Inference on Commodity GPU Servers#155244