CUDA Performance Tuning & Profiling

Performance work is a loop: measure, find the bottleneck, change one thing, and measure again. This lesson uses the three standard measuring tools on a Tesla T4 with CUDA 13.0: CUDA events for timing, Nsight Systems (nsys) for where the time goes across the whole program, and Nsight Compute (ncu) for what a single kernel is doing. Timings moved between runs, so the text gives the range across all the runs.

The Cost of a Kernel Launch

The same 16.7 million additions run once as one launch, then again as 4096 launches of 4096 threads each.

// launch_overhead.cu - the same 16M additions as 1 launch or as 4096 small ones
#include <cstdio>
#include <cuda_runtime.h>

__global__ void add_one(float* x, int offset, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) x[offset + i] += 1.0f;
}

int main() {
  const int N = 1 << 24, PIECES = 4096, P = N / PIECES;   // 4096 elements per piece
  float* x;
  cudaMalloc(&x, N * sizeof(float));
  cudaMemset(x, 0, N * sizeof(float));
  cudaEvent_t a, b, c;
  cudaEventCreate(&a); cudaEventCreate(&b); cudaEventCreate(&c);

  add_one<<<N / 256, 256>>>(x, 0, N);                    // warm-up
  cudaEventRecord(a);
  add_one<<<N / 256, 256>>>(x, 0, N);                    // one launch over everything
  cudaEventRecord(b);
  for (int p = 0; p < PIECES; p++)                       // the same work, piece by piece
    add_one<<<P / 256, 256>>>(x, p * P, P);
  cudaEventRecord(c);
  cudaEventSynchronize(c);

  float one, many, last;
  cudaEventElapsedTime(&one, a, b);
  cudaEventElapsedTime(&many, b, c);
  cudaMemcpy(&last, x + N - 1, sizeof(float), cudaMemcpyDeviceToHost);
  printf("1 launch of %d threads:   %6.3f ms\n", N, one);
  printf("%d launches of %d threads: %6.3f ms  (%.1f us per launch)\n", PIECES, P, many, many * 1000 / PIECES);
  printf("x[N-1] = %.0f (warm-up + big launch + its piece)\n", last);
  cudaFree(x);
  return 0;
}

$ nvcc -O2 -arch=sm_75 -o launch_overhead launch_overhead.cu && ./launch_overhead
1 launch of 16777216 threads:    0.572 ms
4096 launches of 4096 threads: 14.375 ms  (3.5 us per launch)
x[N-1] = 3 (warm-up + big launch + its piece)

Across four runs the single launch took 0.53 to 0.57 ms and the 4096 small ones 12.4 to 15.5 ms, 3.0 to 3.8 µs per launch. The work is identical, and x[N-1] = 3 confirms every element was added three times. Splitting it into small kernels made it 23 to 29 times slower.

Where the Time Went: nsys

nsys profile records the run, and nsys stats turns the recording into tables. The grep keeps only the launch call and the kernel rows.

$ nsys profile -o launches --force-overwrite true ./launch_overhead > /dev/null 2>&1; nsys stats -q --force-export=true --report cuda_api_sum --report cuda_gpu_kern_sum launches.nsys-rep | sed 's/ *$//' | grep -E 'Summary|Name|-----|cudaLaunchKernel|add_one'
 ** CUDA API Summary (cuda_api_sum):
 Time (%)  Total Time (ns)  Num Calls    Avg (ns)       Med (ns)      Min (ns)     Max (ns)    StdDev (ns)           Name
 --------  ---------------  ---------  -------------  -------------  -----------  -----------  -----------  ----------------------
      8.9       17,229,006      4,098        4,204.2        3,938.5        3,259      258,093      4,218.0  cudaLaunchKernel
      0.2          423,192      4,098          103.3          102.0           83        4,859         75.3  cuKernelGetName
 ** CUDA GPU Kernel Summary (cuda_gpu_kern_sum):
 Time (%)  Total Time (ns)  Instances  Avg (ns)  Med (ns)  Min (ns)  Max (ns)  StdDev (ns)             Name
 --------  ---------------  ---------  --------  --------  --------  --------  -----------  --------------------------
    100.0       12,994,588      4,098   3,171.0   2,848.0     2,688   624,050     13,717.1  add_one(float *, int, int)

There are two separate costs. On the CPU, the 4,098 cudaLaunchKernel calls (4096 small, the big one and the warm-up) averaged 4.1 to 4.3 µs each across four profiles, 17 to 18 ms in total. On the GPU, the median add_one took 2.2 to 3.0 µs. That median is a small kernel, because 4,096 of the 4,098 instances are small, and it’s far longer than its 4096 additions need: at 250 GB/s, 32 KB of reads and writes take about 0.13 µs. Each kernel pays a fixed cost to start and finish on the GPU, and here that fixed cost is most of its run time. (cuKernelGetName is the profiler looking up kernel names; the Name filter catches it too.)

The fix is fewer, larger kernels. Practice 2 tries a CUDA graph, which removes most of the CPU-side cost, but not the GPU-side one.

Choosing a Block Size: the Occupancy API

Occupancy is the fraction of an SM’s warp slots that are filled. The CUDA runtime can work it out for a given kernel and block size.

// occupancy.cu - what the occupancy API suggests, and how many blocks fit per SM
#include <cstdio>
#include <cuda_runtime.h>

__global__ void scale(const float* x, float* y, int N) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < N) y[i] = x[i] * 1.1f;
}

int main() {
  int minGrid = 0, block = 0;
  cudaOccupancyMaxPotentialBlockSize(&minGrid, &block, scale, 0, 0);
  printf("suggested block size %d, smallest grid that fills the GPU %d blocks\n", block, minGrid);

  cudaDeviceProp p;
  cudaGetDeviceProperties(&p, 0);
  printf("%d SMs, each up to %d threads and %d blocks\n",
         p.multiProcessorCount, p.maxThreadsPerMultiProcessor, p.maxBlocksPerMultiProcessor);
  for (int bs : {32, 64, 128, 256, 512, 1024}) {
    int blocks = 0;
    cudaOccupancyMaxActiveBlocksPerMultiprocessor(&blocks, scale, bs, 0);
    printf("block %4d: %2d blocks per SM = %4d threads, occupancy %3d%%\n",
           bs, blocks, blocks * bs, 100 * blocks * bs / p.maxThreadsPerMultiProcessor);
  }
  return 0;
}

$ nvcc -O2 -arch=sm_75 -o occupancy occupancy.cu && ./occupancy
suggested block size 1024, smallest grid that fills the GPU 40 blocks
40 SMs, each up to 1024 threads and 16 blocks
block   32: 16 blocks per SM =  512 threads, occupancy  50%
block   64: 16 blocks per SM = 1024 threads, occupancy 100%
block  128:  8 blocks per SM = 1024 threads, occupancy 100%
block  256:  4 blocks per SM = 1024 threads, occupancy 100%
block  512:  2 blocks per SM = 1024 threads, occupancy 100%
block 1024:  1 blocks per SM = 1024 threads, occupancy 100%

A T4 SM holds at most 1024 threads and 16 blocks. With 32-thread blocks the 16-block limit is hit first, at 512 threads, so occupancy can’t pass 50%. Every size from 64 up can fill the SM. cudaOccupancyMaxPotentialBlockSize suggested 1024, the largest block that still reaches full occupancy, and a minimum grid of 40 blocks, one per SM. The API reports what the limits allow, not what’s fastest, so measure before you trust the suggestion.

Theoretical vs Achieved Occupancy: ncu

ncu measures what really happened. sm__maximum_warps_per_active_cycle_pct is the theoretical occupancy and sm__warps_active the achieved one, the average share of warp slots that were really occupied.

// occ_run.cu - the same kernel with 32-thread and 256-thread blocks, for ncu to measure
#include <cstdio>
#include <cuda_runtime.h>

__global__ void scale(const float* x, float* y, int N) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < N) y[i] = x[i] * 1.1f;
}

int main() {
  const int N = 1 << 24;
  float *x, *y;
  cudaMalloc(&x, N * sizeof(float));
  cudaMalloc(&y, N * sizeof(float));
  cudaMemset(x, 0, N * sizeof(float));
  scale<<<N / 32, 32>>>(x, y, N);                 // 32-thread blocks: 1 warp each
  scale<<<N / 256, 256>>>(x, y, N);               // 256-thread blocks: 8 warps each
  cudaDeviceSynchronize();
  printf("launched %d blocks of 32, then %d blocks of 256\n", N / 32, N / 256);
  cudaFree(x); cudaFree(y);
  return 0;
}

$ nvcc -O2 -arch=sm_75 -o occ_run occ_run.cu && ncu --metrics gpu__time_duration.sum,sm__maximum_warps_per_active_cycle_pct,sm__warps_active.avg.pct_of_peak_sustained_active ./occ_run | grep -E '^  scale|duration|warps'
  scale(const float *, float *, int) (524288, 1, 1)x(32, 1, 1), Context 1, Stream 7, Device 0, CC 7.5
    gpu__time_duration.sum                                     ms         1.64
    sm__maximum_warps_per_active_cycle_pct                      %           50
    sm__warps_active.avg.pct_of_peak_sustained_active           %        31.91
  scale(const float *, float *, int) (65536, 1, 1)x(256, 1, 1), Context 1, Stream 7, Device 0, CC 7.5
    gpu__time_duration.sum                                     us       560.51
    sm__maximum_warps_per_active_cycle_pct                      %          100
    sm__warps_active.avg.pct_of_peak_sustained_active           %        77.83

The theoretical numbers match the occupancy API: 50% for 32-thread blocks and 100% for 256. Achieved occupancy is lower in both cases, 31.9% and 77.8% (four runs agreed within 0.1 points), because blocks start and finish at different times and slots sit empty in between. With 32-thread blocks the kernel took 1.64 ms against 0.56 ms for 256, 2.9 times slower for the same work.

A Profiling Workflow

1. Time the whole thing with CUDA events, after a warm-up launch, averaged over several runs.
2. Run nsys to see where the time goes: kernels, copies, launch calls or gaps. Fix the biggest item first.
3. Run ncu on the kernels that dominate. Compare the numbers with what the code should do: bytes from DRAM against the data size, achieved against theoretical occupancy.
4. Change one thing, such as the block size, fusing kernels or the memory layout, and re-measure. The roofline lesson shows kernel fusion measured the same way.

Common Gotchas

• Timing without a warm-up. The first launch includes one-time setup. Time later launches.
• Many tiny kernels. Each one costs microseconds on both the CPU and the GPU. Do more work per launch.
• Trusting the occupancy API’s suggestion. It reports what fits, not what’s fastest. 100% theoretical occupancy still left the kernel at 78% achieved.
• Profiling changes timing. ncu replays kernels and locks clocks, so its durations aren’t directly comparable with event timings from a normal run.

Practice

1. Split the work into 256 launches instead of 4096. Does the per-launch cost change?

$ nvcc -O2 -arch=sm_75 -o p1_launch256 p1_launch256.cu && ./p1_launch256
1 launch of 16777216 threads:    0.534 ms
256 launches of 65536 threads:  1.241 ms  (4.8 us per launch)
x[N-1] = 3 (warm-up + big launch + its piece)

2. Record the 4096 small launches once as a CUDA graph and replay the graph with one cudaGraphLaunch. How much faster is the replay?

// p2_graph.cu - record the 4096 small launches once as a CUDA graph, then replay it
#include <cstdio>
#include <cuda_runtime.h>

__global__ void add_one(float* x, int offset, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) x[offset + i] += 1.0f;
}

int main() {
  const int N = 1 << 24, PIECES = 4096, P = N / PIECES;
  float* x;
  cudaMalloc(&x, N * sizeof(float));
  cudaMemset(x, 0, N * sizeof(float));
  cudaStream_t s;
  cudaStreamCreate(&s);

  cudaGraph_t g;
  cudaGraphExec_t ge;
  cudaStreamBeginCapture(s, cudaStreamCaptureModeGlobal);  // record, don't run
  for (int p = 0; p < PIECES; p++) add_one<<<P / 256, 256, 0, s>>>(x, p * P, P);
  cudaStreamEndCapture(s, &g);
  cudaGraphInstantiate(&ge, g, 0);
  cudaGraphLaunch(ge, s);                                   // warm-up replay
  cudaStreamSynchronize(s);

  cudaEvent_t a, b, c;
  cudaEventCreate(&a); cudaEventCreate(&b); cudaEventCreate(&c);
  cudaEventRecord(a, s);
  for (int p = 0; p < PIECES; p++) add_one<<<P / 256, 256, 0, s>>>(x, p * P, P);
  cudaEventRecord(b, s);
  cudaGraphLaunch(ge, s);                                   // all 4096 kernels in one call
  cudaEventRecord(c, s);
  cudaEventSynchronize(c);

  float plain, graph, last;
  cudaEventElapsedTime(&plain, a, b);
  cudaEventElapsedTime(&graph, b, c);
  cudaMemcpy(&last, x + N - 1, sizeof(float), cudaMemcpyDeviceToHost);
  printf("4096 separate launches: %6.3f ms\n", plain);
  printf("one graph launch:       %6.3f ms\n", graph);
  printf("x[N-1] = %.0f\n", last);
  return 0;
}

$ nvcc -O2 -arch=sm_75 -o p2_graph p2_graph.cu && ./p2_graph
4096 separate launches: 12.221 ms
one graph launch:        8.087 ms
x[N-1] = 3

3. Profile the kernel with 64-thread blocks. The occupancy API says 64 already reaches 100%. Is it as fast as 256?

$ nvcc -O2 -arch=sm_75 -o p3_occ64 p3_occ64.cu && ncu --metrics gpu__time_duration.sum,sm__maximum_warps_per_active_cycle_pct,sm__warps_active.avg.pct_of_peak_sustained_active ./p3_occ64 | grep -E '^  scale|duration|warps'
  scale(const float *, float *, int) (262144, 1, 1)x(64, 1, 1), Context 1, Stream 7, Device 0, CC 7.5
    gpu__time_duration.sum                                     us       678.56
    sm__maximum_warps_per_active_cycle_pct                      %          100
    sm__warps_active.avg.pct_of_peak_sustained_active           %        62.05
  scale(const float *, float *, int) (65536, 1, 1)x(256, 1, 1), Context 1, Stream 7, Device 0, CC 7.5
    gpu__time_duration.sum                                     us       560.32
    sm__maximum_warps_per_active_cycle_pct                      %          100
    sm__warps_active.avg.pct_of_peak_sustained_active           %        77.81

4. Instead of one thread per element, launch a fixed grid of 160 blocks (4 per SM, so 1024 threads per SM) in which each thread loops over the array with a grid-stride loop. Is it faster?

// p4_gridstride.cu - one thread per element vs a fixed grid of 160 blocks that loops
#include <cstdio>
#include <cuda_runtime.h>

__global__ void add_one(float* x, int n) {                 // one element per thread
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) x[i] += 1.0f;
}

__global__ void add_one_stride(float* x, int n) {          // each thread walks the array
  for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n; i += gridDim.x * blockDim.x)
    x[i] += 1.0f;
}

int main() {
  const int N = 1 << 24;
  float* x;
  cudaMalloc(&x, N * sizeof(float));
  cudaMemset(x, 0, N * sizeof(float));
  cudaEvent_t a, b, c;
  cudaEventCreate(&a); cudaEventCreate(&b); cudaEventCreate(&c);
  add_one<<<N / 256, 256>>>(x, N);                         // warm-up both
  add_one_stride<<<160, 256>>>(x, N);
  cudaEventRecord(a);
  for (int r = 0; r < 10; r++) add_one<<<N / 256, 256>>>(x, N);
  cudaEventRecord(b);
  for (int r = 0; r < 10; r++) add_one_stride<<<160, 256>>>(x, N);   // 4 blocks per SM
  cudaEventRecord(c);
  cudaEventSynchronize(c);
  float one, stride, last;
  cudaEventElapsedTime(&one, a, b);
  cudaEventElapsedTime(&stride, b, c);
  cudaMemcpy(&last, x + N - 1, sizeof(float), cudaMemcpyDeviceToHost);
  printf("65536 blocks, 1 element per thread: %.3f ms\n", one / 10);
  printf("160 blocks, grid-stride loop:       %.3f ms\n", stride / 10);
  printf("x[N-1] = %.0f\n", last);
  cudaFree(x);
  return 0;
}

$ nvcc -O2 -arch=sm_75 -o p4_gridstride p4_gridstride.cu && ./p4_gridstride
65536 blocks, 1 element per thread: 0.531 ms
160 blocks, grid-stride loop:       0.646 ms
x[N-1] = 22