CUDA kernel fusion and what a launch really costs
Day 20 fused a Sobel kernel and a threshold kernel into one, and the time fell from 2.565 ms to 1.323 ms on a Tesla T4. Almost everyone reads that number the same way: two launches became one, so a launch must be expensive.
The launch count does not explain that result. The same table says the two-kernel version moved 218 MB and the fused one moved 84 MB. The speedup was 1.9 and the traffic ratio was 2.6.
Nothing in that row was launch cost, and nothing had to be: on that card a launch is microseconds and those kernels ran for milliseconds.
Fusion can reduce memory traffic and increase register use. This lesson measures both effects on a four-stage elementwise chain. It finds the sizes where each cost affects run time.
Two costs, and they are not in the same unit
Fusing a chain deletes two different things.
The first is launch overhead: a fixed cost to get a kernel started and finished, paid whatever the kernel does. It is microseconds, and it does not care how big your arrays are. Four launches cost four times one launch at every size.
The second is a round trip through global memory. Every kernel in a chain reads its input from memory and writes its output back, so a value that only the next kernel will ever look at is still stored, usually evicted, and fetched again. That cost is bytes, and bytes scale with the problem.
Here is the chain. Four stages, the tail of a transformer's feed-forward block written the way a first version writes it:
applyScaleBias<<<blocks, 256>>>(d_x, d_t1, n); // t1 = x * g + b
applyGelu<<<blocks, 256>>>(d_t1, d_t2, n); // t2 = gelu(t1)
applyAddResidual<<<blocks, 256>>>(d_t2, d_r, d_t3, n);// t3 = t2 + r
applyClampRange<<<blocks, 256>>>(d_t3, d_y, n); // y = clamp(t3)
Count the floats. Stages 1 and 2 each read one and write one. Stage 3 reads two, because the residual is a second input, and writes one.
Stage 4 reads one and writes one. The staged chain moves nine floats, or 36
bytes, per element. The fused version reads x and r, then writes y, for
three floats or 12 bytes.
Write that count down before you measure anything, because it is the prediction. The traffic ratio is three and the launch ratio is four, and those are different numbers, so the measurement can say which one you were the timing reflects. A wrong count can still appear to fit a speedup: a chain of four identical two-float stages would move 32 bytes, and a reader with that model would still see a speedup and conclude the model was right.
Diagram: one chain, three fusion depths. Three horizontal bands. Each band has kernel boxes left to right along a time axis, with the global memory strip drawn underneath and one arrow per float read or written between a box and the strip. Band 1, four kernels: four boxes, nine arrows, three intermediate buffers drawn on the strip between them. Caption "4 launches, 36 bytes per element." Band 2, fused in pairs: two boxes, five arrows, one intermediate buffer left on the strip. Caption "2 launches, 20 bytes per element." Band 3, fused: one box, three arrows, no intermediate on the strip at all; the two values that were buffers are now drawn as registers attached to the box. Caption "1 launch, 12 bytes per element." Alt text: "One elementwise chain at three fusion depths. Four kernels move thirty-six bytes per element, two move twenty, one moves twelve, and the launches fall from four to one."
The intuition that fusion is about launch count
The belief comes from somewhere real. Machine learning frameworks fuse elementwise operations and often describe the gain as lower launch overhead. For a model that issues hundreds of small pointwise kernels on tensors of a few thousand elements, launch overhead can dominate.
Day 56 measures that case.
Memory traffic dominates once a kernel runs much longer than its launch cost. Check the units before accepting a claim. Sixteen million elements at 36 bytes each is 604 MB.
At the measured copy rate, the staged chain spends milliseconds moving data. Four launches take microseconds, three orders of magnitude less.
Fusion also has a cost. Each intermediate removed from memory may need a register, and an SM has a fixed register file. More registers per thread can reduce occupancy and leave fewer warps to hide memory latency.
Day 17 measured that trade on one kernel, and day 45 studies it directly. Here it is a possible cost of fusion.
Holding the arithmetic still and moving the memory
Full program in
code/day48-fusion/fusion.cu. Four
properties make the timing mean something.
Both versions call the same four functions. The stages are ordinary
__device__ functions, so the staged kernels and every fused kernel run the
same arithmetic in the same order. Anything the clock finds between them is
traffic and launches, not maths.
The byte counts come from the constants and print above the times. The program cannot compute a GB/s figure that disagrees with the model it is being judged against, because both come from the same two numbers.
The copy baseline is measured here. A copyFloor kernel reads the
buffer and writes it back with no arithmetic. No elementwise chain over the
same data can beat it, and measuring it in this process keeps this program's
block shape, buffer size and clock state out of the comparison.
Events, and a warm-up per kernel. A host clock around a launch measures the launch, which day 9 takes apart, and the first launch of each kernel includes its own module load under lazy loading.
The fused chain takes nine lines:
__global__ void chainFused(const float* __restrict__ x,
const float* __restrict__ r, float* __restrict__ out,
size_t n) {
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
if (i < n) {
float v = scaleBias(x[i]);
v = gelu(v);
v = addResidual(v, r[i]);
out[i] = clampRange(v);
}
}
chainFusedWide<K> tests the register cost. It is the same
chain with each thread carrying K elements instead of one: same total work,
same total bytes, same arithmetic in the same order. The three loops are
separate so that all K values are live at once.
It puts K memory requests in flight per thread and makes the compiler find
K registers to keep them in.
template <int kWide>
__global__ void chainFusedWide(const float* __restrict__ x,
const float* __restrict__ r,
float* __restrict__ out, size_t n) {
const size_t base =
blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
const size_t stride = gridDim.x * static_cast<size_t>(blockDim.x);
float v[kWide];
#pragma unroll
for (int k = 0; k < kWide; ++k) {
const size_t i = base + static_cast<size_t>(k) * stride;
v[k] = (i < n) ? scaleBias(x[i]) : 0.0f;
}
#pragma unroll
for (int k = 0; k < kWide; ++k) {
const size_t i = base + static_cast<size_t>(k) * stride;
v[k] = addResidual(gelu(v[k]), (i < n) ? r[i] : 0.0f);
}
#pragma unroll
for (int k = 0; k < kWide; ++k) {
const size_t i = base + static_cast<size_t>(k) * stride;
if (i < n) {
out[i] = clampRange(v[k]);
}
}
}
The elements a thread owns are one whole grid apart, not adjacent, so for
each k the 32 lanes of a warp still read 32 consecutive
floats and every load coalesces. That is the
grid-stride loop from day 8 with its trip count
known at compile time.
Note.
chainFusedWide<1>should be the same kernel aschainFused: one element, one guard, no array. The program times both anyway. If those two rows disagree, the scaffolding around the template is costing something and every other row in that table is measuring it too.
Registers and occupancy come from cudaFuncGetAttributes and
cudaOccupancyMaxActiveBlocksPerMultiprocessor, which are runtime API calls.
These calls need no performance counters, profiler, or root access. Learners
can reproduce the second table without access to hardware counters.
Results
Re-verified on a Tesla T4 with driver 580.173.02 and CUDA 13.0 V13.0.88 on 2026-09-02. Correctness, the 3.00 byte ratio and the large-size 3.09x fusion speedup reproduced. Compiler changes removed the old width/occupancy drop: CUDA 13 used at most 56 registers for every width, kept 100 percent theoretical occupancy and showed no spill.
The original CUDA 12.6 table remains below; the CUDA 13 transcript and Nsight Systems report are listed in front matter.
Originally measured on a Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with
nvcc -std=c++17 -O3 -arch=sm_75, timed with CUDA events, mean of 10 runs
after 3 warm-ups at the 16,777,216-element timed size. Captured 2026-09-01
on the project's verification node; the full four-part transcript is in
the page's evidence file.
Part 1: what one launch costs, 1000 back-to-back launches per timed run
launch shape per launch (us)
------------------------------------- ---------------
doNothing, 1 block of 1 thread 2.668
doNothing, the chain's own grid 71.668
| Part 2 row | ms | bytes/elem | GB/s | x copy |
|---|---|---|---|---|
copyFloor |
0.548 | 8 | 244.8 | 1.00 |
| four stages, summed | 2.429 | 36 | 248.6 | 1.02 |
| staged chain, timed as one | 2.419 | 36 | 249.7 | 1.02 |
chainFused |
0.785 | 12 | 256.5 | 1.05 |
Part 3: the same two chains, swept over n
n staged (ms) fused (ms) staged/fused staged us/launch
---------- ----------- ----------- ------------ ----------------
1024 0.0162 0.0058 2.79 4.050
16384 0.0107 0.0045 2.38 2.686
262144 0.0299 0.0051 5.90 7.475
4194304 0.6128 0.1984 3.09 153.194
16777216 2.4264 0.7864 3.09 606.606
Part 4: chainFused widened, same work and same bytes every row
per thread blocks regs spill B blocks/SM occupancy ms GB/s
---------- ------- ----- ------- --------- --------- ------- --------
1 65536 11 0 4 100.0% 0.792 254.2
2 32768 18 0 4 100.0% 0.798 252.3
4 16384 23 0 4 100.0% 0.805 250.1
8 8192 41 0 4 100.0% 0.817 246.6
16 4096 80 0 3 75.0% 0.832 242.1
32 2048 174 0 1 25.0% 1.191 169.1
The four claims, against the card:
The first prediction held. Staged over fused is 3.08 at the timed size against a byte ratio of 3.00 and a launch ratio of 4. Every row runs within 5 percent of the copy floor, so the chain is memory bound at an arithmetic intensity of a few flops per byte, and the speedup matches the lower byte count.
The time-per-launch prediction held, but the ratio did not. The staged chain's time per launch falls to 2.686 us at 16,384 elements, on top of part 1's 2.668 us empty-kernel figure, so the last column converged as predicted. The staged-to-fused ratio did not climb toward four; it fell to 2.38, because the fused version's single launch includes the same fixed costs, and shrinking work amplifies them on both sides of the fraction. The 262,144 row shows that microsecond-scale timings vary: its 5.90 came with a per-launch figure above the rows either side of it.
The width sweep still shows no gain, but the compiler reason changed. Under CUDA 12.6, registers grew from 11 to 174 and occupancy fell to 25 percent at width 32 without spilling. Under CUDA 13.0, the same source used 13, 23, 34, 55, 56 and 56 registers across the sweep: every width kept 100 percent theoretical occupancy and zero spill bytes.
The narrowest width was still fastest, 0.784 ms against 0.873 at width 32. Widening did not help, but the new artifact disproves a stable link between width and register use. That relationship belongs to a particular compiler build and must be measured rather than carried forward.
The bitwise-equality prediction failed on 1,850,839 elements. The two chains differ on 11 percent of outputs, by at most 2.38e-07. The prediction hedged for exactly this, and the likely cause is the compiler contracting multiplies and adds into fused multiply-adds differently across the fused and staged bodies, which day 47 measures. The program counted the differences instead of assuming equality.
The lesson includes two types of profiler artifact. The two halves of the question need different tools.
content/profiles/kernel-fusion/kernel-fusion.nsys-rep is the
Nsight Systems timeline: four kernel bars with
three gaps between them, then one bar. The gaps show launch time at the
timeline's scale. Nsight Systems does not need performance counters, so it runs as a
normal user and it runs on Colab.
content/profiles/kernel-fusion/kernel-fusion-<kernel>.ncu-rep, one per
kernel, is the Nsight Compute side, and it needs
performance counters. Plain ncu fails on a stock driver. It returns
ERR_NVGPUCTRPERM, in full: The user running <tool_name/application_name> does not have permission to access NVIDIA GPU Performance Counters or the Hardware Event System on the target device
(https://developer.nvidia.com/nvidia-development-tools-solutions-err_nvgpuctrperm-permission-issue-performance-counters
, checked 2026-08-30).
The cause is the driver parameter
RmProfilingAdminOnly, which is 1 by default in
/proc/driver/nvidia/params. sudo ncu works where you have root.
This project verified that ordinary-user NCU is blocked on the measured node when that parameter is 1. It has not verified current free-Colab behavior, so that question remains open.
The shipped reports provide a path either way. The two
metrics used here are dram__bytes_read.sum and
dram__bytes_write.sum, in the MemoryWorkloadAnalysis section, and the
exercise below is doable from the text export alone.
| Kernel | dram__bytes_read.sum |
dram__bytes_write.sum |
bytes/elem |
|---|---|---|---|
applyScaleBias |
73.380320 MB | 73.805600 MB | 8.8 |
applyGelu |
74.462368 MB | 74.150336 MB | 8.9 |
applyAddResidual |
145.788576 MB | 74.895648 MB | 13.2 |
applyClampRange |
73.378816 MB | 73.939456 MB | 8.8 |
chainFused |
147.114048 MB | 75.105728 MB | 13.2 |
Each row sits a few percent over its algorithmic 8 or 12 bytes per element because the profiled pass runs cold, with the caches flushed. The pattern matters: the two rows that read two buffers move twice the read traffic, and the fused chain's total is a third of the staged chain's four rows summed.
The program counts required traffic. The profiler measures DRAM traffic.
Day 30 could not compare those counts because
that requires the profiler's dram__bytes metric. This lesson compares them.
Run it yourself
Run the program on a supported CUDA GPU. The README uses this command for the shipped target:
nvcc -std=c++17 -O3 -arch=sm_75 -o fusion fusion.cu
The full run does not fit Compiler Explorer's time limit. Six 64 MiB buffers and around thirty thousand launches across the four parts use up Compiler Explorer's 20 seconds before the compile is counted, and the compile alone is slow because the widest fused kernel unrolls 32 copies of a GELU.
Shrinking it to fit would change the traffic under test. If you have no GPU at all, read /setup/learn-cuda-without-a-gpu.
Exercise
Fuse the chain half way: one kernel for stages 1 and 2, a second for stages 3 and 4, with one intermediate buffer between them. Work out its bytes per element from the code before you compile it, predict both ratios, then measure and say which prediction held.
Time: 30 to 45 minutes. Submit: the two kernels, the row you added to the part 2 table, and your two predicted ratios beside the measured ones.
Check: the program compares every version against the double reference at 1,049,187 elements, an odd size that no power-of-two launch shape covers exactly, so every bounds check fires on every launch. It prints the first bad index with both values, so adding your pair to that comparison checks correctness.
Two mistakes account for most failures here: an intermediate buffer that is one of the three the staged chain already uses, so your second kernel reads what the staged run left there, and a second kernel that re-applies stage 2. Then add your row to the timed table with its own byte count, because two rows moving different byte counts are not comparable in milliseconds.
Hint 1
You are not deleting a launch here. You are deleting a store, and the load that follows it. Which of the three intermediate buffers stops existing, and which one does not?
Hint 2
Write down the floats each of your two kernels reads and writes per element, and do not forget that one of them has a second input. The staged chain is nine floats, and the fused chain is three.
What is yours, and what do those three numbers predict for the two ratios you were asked for?
Solution
Five floats, so 20 bytes per element. The first kernel reads x and writes
the one surviving intermediate; the second reads that intermediate, reads the
residual, and writes y. The predicted ratios are 36 over 20, which is 1.8
against the staged chain, and 20 over 12, which is 1.67 against the fused
one.
Both predictions hold only while the chain is bound by traffic. At the small end of the part 3 sweep they both approach the launch ratios, 4 over 2 and 2 over 1, because at that size neither version is moving enough bytes for bytes to matter.
Estimate fusion from the traffic it removes. Count the floats each version reads and writes per element. If that count does not change, fusing will remove one launch and nothing else. A launch takes microseconds.
Pitfalls
You fused two kernels and nothing got faster. Count the bytes before you attribute the result to the GPU. If the intermediate was small enough to fit in L2, the second kernel was already reading a cache hit and there was no round trip to delete, so all fusion removed was a launch. Day 34 hit the same shape from the other side: a shared-memory tile that cut global reads per cell from five to 1.328 and made the kernel slower, because the reads it removed were cache hits and not DRAM traffic.
Your fused kernel is slower than the four it replaced. Check its register
use. Read numRegs out of cudaFuncGetAttributes for both, and
cudaOccupancyMaxActiveBlocksPerMultiprocessor for both before using
a profiler. Neither call needs performance counters.
If the fused kernel needs more registers per thread than the GPU allows at full occupancy, you have traded latency hiding for traffic, and whether that was a good trade is a measurement. Day 45 is the lesson.
You timed the four stages separately and added them up. Four kernels measured in four loops is not one chain measured in one region. Time the sequence you actually ship, and keep the per-stage numbers as a breakdown rather than as the answer.
Your fused answer differs from the staged answer in the last bits. An intermediate that goes to memory is rounded to float on the way; one that stays in a register may be carried into a fused multiply-add the staged version could not use. Day 68 covers bitwise reproducibility.
Report the difference, do not assert it away.
You used CUDA graphs to fix a traffic problem. A graph collapses the cost of submitting a fixed sequence of launches. It moves no bytes and deletes no round trips, so on the chain above it can remove the microseconds but not the milliseconds. Day 56 captures this exact chain and measures what is left.
You profiled as a normal user and got nothing. ncu reports
ERR_NVGPUCTRPERM. On a machine where you have root, sudo ncu works when
RmProfilingAdminOnly is 1. The measured node proves that local restriction,
not current free-Colab behavior; the hosted-tier question remains open.
The page ships reports so the exercise never depends on counter access.
Go deeper
- CUDA Programming Guide, "Lazy Loading", for why the first launch of each kernel costs more than the rest and what that does to a benchmark without a warm-up: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/lazy-loading.html (checked 2026-08-30)
- CUDA C++ Best Practices Guide, "Occupancy" and "Register Pressure", for the occupancy API used in part 4 and the register trade it reports: https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html (checked 2026-08-30)
- Nsight Compute CLI reference, for
--set,--kernel-name,--launch-skipand--clock-control, which are the four options the capture in the README uses: https://docs.nvidia.com/nsight-compute/NsightComputeCli/index.html (checked 2026-08-30) cuda-samples,cpp/3_CUDA_Features/simpleCudaGraphs, the same idea from the launch-overhead side, and the sample day 56 builds on: https://github.com/NVIDIA/cuda-samples/tree/master/cpp/3_CUDA_Features/simpleCudaGraphs (checked 2026-08-30)- Programming Massively Parallel Processors, 4th edition, chapter 6, on performance considerations and thread coarsening, which is what widening the fused kernel does: https://shop.elsevier.com/books/programming-massively-parallel-processors/hwu/978-0-323-91231-0
Next
Day 49 builds a roofline for a measured GPU and puts this chain on it, where the staged and fused versions are two points at different arithmetic intensities doing identical work. Day 56 takes this exact chain, captures it as a CUDA graph, and measures how much of part 1's number survives when the launches stop being launches. Both sit in module 5, and the GPU kernel engineer hub is where the rest of it goes.