Kernel Fusion and the Cost of a Launch Boundary

Two kernels that each read an array, transform it, and write it back move twice the data that one kernel doing both transformations would. Whether that saving is worth taking depends on three quantities: the fixed cost of the launch itself, the traffic eliminated at the intermediate boundary, and the cost of the fusion in registers, occupancy, and scheduling.

What a launch actually costs

A launch is not free but it is not the dominant cost people assume. The device-side work of scheduling blocks is small; the expensive parts are the host-side submission and, more importantly, the implicit serialization at the boundary. A kernel must complete before a dependent kernel’s blocks can read its output, so the second kernel cannot start until the first has drained. That drain is not instantaneous: the last wave of blocks finishes over a period that includes the block granularity and the memory writes still in flight.

The serialization has two consequences. The first is idle time at the boundary: the machine empties at the end of kernel one and refills at the start of kernel two, and during both transitions the device is partially occupied. The second is that any tail effect in the first kernel, described in wave quantization, is paid in full and cannot be overlapped with work from the second.

On a stream that allows concurrent kernels, two independent kernels can overlap and the boundary cost is mostly hidden. For a dependent pair, the second kernel’s blocks cannot read the first kernel’s output reliably until the first has completed, so overlap is limited to the extent the dependency is finer-grained than the kernel boundary. On hardware that supports fine-grained dependencies between blocks, some overlap is possible; without that, the boundary is a full barrier.

The fixed per-launch cost on the host side is on the order of a few microseconds of CPU time. For a kernel that runs for milliseconds, that is negligible. For a kernel that runs for tens of microseconds, it is a large fraction, and a sequence of such kernels is dominated by launch overhead rather than by computation. Whether that is the case is determined by measuring the kernel duration, not by assuming.

The traffic argument

The stronger argument for fusion is traffic. Consider two elementwise kernels applied in sequence. The first reads $n$ elements and writes $n$. The second reads $n$ and writes $n$. Total DRAM traffic is $4n$ elements’ worth, of which the $2n$ at the boundary is pure overhead: the intermediate array is written and immediately read, and never observed by anything else.

Fusing them into one kernel removes the intermediate array entirely, leaving $2n$ of traffic. The arithmetic is unchanged. If both kernels were bandwidth bound, which elementwise kernels almost always are, the fused kernel runs in half the time. This is the clearest possible case and the one that justifies fusion as a default.

// Two passes: 4n element transfers, two launches.
scale<<<grid, block>>>(n, a, x, tmp);     // read x, write tmp
bias <<<grid, block>>>(n, b, tmp, y);     // read tmp, write y

// One pass: 2n element transfers, one launch.
fused<<<grid, block>>>(n, a, b, x, y);    // read x, write y

The traffic saving is not always a factor of two. If the second kernel is compute bound and the first is bandwidth bound, the fused kernel’s arithmetic cost dominates and the traffic saving is a smaller fraction of the total. The position of each kernel on the roofline determines how much the fusion can buy, and a fusion that removes traffic from a compute-bound kernel buys nothing.

The saving also depends on whether the intermediate fits in cache. If the intermediate array is small enough that the second kernel reads it entirely from L2, the boundary traffic never reaches DRAM and the fusion saves cache bandwidth rather than DRAM bandwidth. That is still a saving, but a smaller one, and it is measured in the L2 counters rather than the DRAM ones.

What fusion costs

Fusion is not free, and the costs are the reason not to do it universally.

Register pressure. A fused kernel holds the live values of both stages simultaneously. Where the first kernel needed a few registers and the second needed a few, the fused kernel needs the union plus whatever the intermediate values occupy. If that pushes the kernel over a register threshold, occupancy falls, and the loss can exceed the traffic gain. The calculation is the same one as in occupancy as a latency-hiding budget, run with the fused register count.

Reduced parallelism. Two separate kernels can be launched with grids tuned to their own work. A fused kernel has one grid, and if the stages have different natural parallel decompositions, one of them is served badly. A stage that wants one thread per output row fused with a stage that wants one thread per element forces a compromise in the thread mapping.

Loss of scheduling flexibility. Separate kernels let the scheduler interleave the tail of one with the head of another when there is no dependency. A fused kernel has one wave structure, and the fusion may create a longer critical path.

Code complexity and reuse. A fused kernel is specific to the pair. If either stage is used elsewhere on its own, the fusion duplicates logic or forces a parameterized kernel with a branch, which reintroduces the divergence cost discussed in issue limits and divergence.

Different block sizes and shared memory. Two kernels can each request the shared memory they need. A fused kernel requests the sum or the maximum, and if the sum exceeds what the block can reserve, the fusion is impossible without restructuring.

The shapes that fuse well and the ones that do not

The favourable case is producer-consumer with an elementwise or near-elementwise relationship, where each output element depends on a small, local set of input elements. A stencil, a pointwise activation, a scale-and-bias, a normalization: all of these have the property that the mapping from input to output is statically known and local, so a single kernel can compute each output directly from the inputs without materializing the intermediate.

The unfavourable case is a global dependency. A reduction followed by a normalization needs the whole reduction before any output can be computed, so the two stages cannot be merged without a grid-wide barrier or a multi-pass algorithm. A sort followed by a gather has the same shape. These are genuinely two phases, and the fusion would require either a cooperative launch with grid synchronization or a different algorithm entirely. The intermediate traffic is the price of the phase boundary, and it is often the right price to pay.

Between the two is the case where fusion is possible but requires recomputation. A stage that computes a value used by several downstream outputs can either materialize it once (a separate kernel) or recompute it per output (fused). The trade is traffic against arithmetic. Recomputing is better when the arithmetic is cheap and the traffic is expensive; materializing is better when the arithmetic is expensive and the traffic is cheap. Since the arithmetic is often a transcendental or a division and the traffic is often a full cache line, the decision is not obvious and depends on the reuse factor.

Structure	Fusion verdict	Reason
Elementwise then elementwise	Fuse	Boundary traffic eliminated, no added pressure
Elementwise then reduction	Fuse with care	Reduction needs the whole array; partial sums may be fused
Reduction then elementwise	Do not fuse	Grid-wide dependency
Stencil then stencil	Fuse	Locality preserved, halo recomputed
Sort then gather	Do not fuse	Global permutation
Two stages with different block sizes	Case by case	Grid compromise may cost more than the traffic

Measuring whether the fusion helped

The measurement has to isolate the two effects, because a fusion that helps for the wrong reason will not generalize and may regress on other inputs.

Run the unfused version and record total DRAM bytes and duration. Run the fused version and record the same. The byte difference should match the predicted boundary traffic; if it does not, the intermediate was being cached and the traffic model was wrong. The duration difference divided by the byte difference gives the effective bandwidth of the eliminated traffic, which should be near the practical ceiling from achieved versus peak throughput. If it is far below, the fused kernel has lost occupancy or gained divergence and the traffic saving is being eaten.

The regression case is worth testing explicitly. A fusion that helps at one size may hurt at another, because the intermediate array’s cache residency depends on size. A small intermediate that fits in L2 makes the unfused version cheap, and the fused version’s higher register pressure then makes it slower. The crossover is a function of the array size and the cache capacity, and it should be measured rather than assumed.

Fusion as a scheduling decision, not a code transformation

The framing that makes the decision tractable is to treat fusion as a choice about where the boundary between phases sits. Every boundary costs the traffic of the values crossing it plus the serialization at the boundary. Fusion removes a boundary; it does not remove the dependency, which is why the shapes with global dependencies cannot be fused without changing the algorithm.

A kernel sequence is therefore a design with a cost per boundary, and the question is which boundaries are worth paying for. A boundary between two elementwise stages is pure overhead. A boundary at a global reduction is structural, and the intermediate is the algorithm’s working set. A boundary between stages with different parallelism is a scheduling convenience that costs traffic. Naming which of the three a given boundary is answers whether to fuse it, and the traffic counters confirm the answer.