Kernel Fusion in NVIDIA CUDA: Optimizing Memory Traffic and Launch Overhead | NVIDIA Technical Blog

Technical Blog

Subscribe

Related Resources

Developer Tools & Techniques

English中文

Kernel Fusion in NVIDIA CUDA: Optimizing Memory Traffic and Launch Overhead

Jul 10, 2026

By Daniel Rodriguez

Like

Discuss (0)

There are many ways to optimize code for GPUs. In this post, you’ll learn how kernel fusion can improve memory bandwidth and reduce kernel launch overhead, along with multiple ways to apply it in NVIDIA CUDA code.

A common bottleneck when writing GPU code is that GPU compute is so fast that even high-bandwidth device memory doesn’t use the GPU kernel fully. Kernel fusion addresses this by combining multiple GPU operations into a single device kernel, so intermediate results don’t need to round-trip through global memory or require separate kernel launches.

This post is about fusing kernel bodies so intermediate results stay in registers, and memory transfers are reduced. CUDA Graphs provide another kind of fusion, but at a different layer. A graph captures a sequence of kernel launches, memory copies, and synchronizations into one reusable object that the host can dispatch with a single call. It doesn’t fuse kernel bodies.

Kernel fusion in practice

Kernels inside a graph run separately and pass intermediate results through global memory. Wrapping a naive baseline in a graph shaves microseconds off the host side, but the one GiB round-trip through the intermediate buffer remains unchanged. The two approaches are complementary and can be used together.

Let’s take a simple example: sum(abs(x)). We read an array, take the absolute value of each element, and return the sum.

A naive implementation uses two kernels: one to compute abs into a temporary intermediate buffer of the same size as the input, and another to reduce that buffer to a single number. This works, but it’s slow in a way that shows the broader problem.

template<br>__global__ void abs_kernel(Config config,<br>cuda::std::span in,<br>cuda::std::span tmp)<br>/* Base SIMT implementation using fsabs */

template<br>__global__ void sum_kernel(Config config,<br>cuda::std::span tmp,<br>cuda::std::span out)<br>/* CUB block-reduce + atomic add into out[0] */

int main()<br>/* device, stream, and buffers via cuda::make_buffer<br>* See the CCCL Runtime post for the full setup. */

auto config = cuda::distribute(in.size());

cuda::launch(stream, config, abs_kernel, in, tmp);<br>cuda::launch(stream, config, sum_kernel, tmp, out);

The C++ examples featured throughout this post use the latest NVIDIA CCCL runtime interfaces, including cuda::std::span, cuda::launch, and cuda::make_buffer—which debuted in CUDA 13.2. To explore these primitives further, refer to CCCL Runtime: A Modern C++ Runtime for CUDA.

If we look at an NVIDIA Nsight Systems profile of the full implementation, it looks like this:

Figure 1. Nsight Systems profile of non-fused kernels for sum(abs(x)). Each kernel is executed independently, and an intermediate result is needed

Manual kernel fusion

The first thing we could do is rewrite these two kernels as a single sum_abs_kernel. This removes the intermediate buffer, which exists only to pass data between kernels. If both operations live in the same kernel, the buffer disappears, and the performance of the operation improves.

template<br>__global__ void sum_abs_kernel(Config config,<br>cuda::std::span x,<br>cuda::std::span out)<br>using BlockReduce = cub::BlockReduce;<br>__shared__ typename BlockReduce::TempStorage temp_storage;

// Grid-stride loop. abs() is computed inline -- no tmp buffer.<br>float thread_sum = 0.0f;<br>const auto tid = cuda::gpu_thread.rank(cuda::grid, config);<br>const auto stride = cuda::gpu_thread.count(cuda::grid, config);<br>for (size_t i = tid; i r(out[0]);<br>r.fetch_add(block_sum, cuda::memory_order_relaxed);

int main()<br>/* device, stream, and buffers -- see the CCCL Runtime post. */

// Fixed grid of 1024 blocks for the grid-stride pattern.<br>auto config = cuda::make_config(cuda::make_hierarchy(<br>cuda::grid_dims(1024),<br>cuda::block_dims()));<br>cuda::launch(stream, config, sum_abs_kernel, x, out);

fabsf(x[i]) is computed inline, in a register, instead of being read from an intermediate buffer that another kernel wrote.

A grid-stride loop lets each thread accumulate multiple elements into a single register (thread_sum), so partial sums never touch global memory.

cub::BlockReduce handles the per-block reduction in shared memory, and cuda::atomic_ref lets one thread per block add its partial sum to the global result.

If we look at an Nsight Systems profile of the fused implementation, it collapses accordingly:

Figure 2. Nsight Systems profile of a manually fused kernel. A sum_abs_kernel is a single kernel that does the whole operation

MetricNaiveManual fusionKernels Launched21Intermediate bufferYesNoBytes moved through global memory3 GB1 GBTime3.51 ms = 2.28ms + 1.23 ms1.18 msEffective memory bandwidth855 GiB/s851 GiB/sSpeedup vs. naive–3 timesTable 1....