CUDA Performance Tuning & Profiling

Advanced Mindset: Performance is Feedback (Theory)

CUDA performance tuning is iterative:

1. Measure (profile)
2. Diagnose (what stalls?)
3. Change (tiling, vectorization, fusion, launch config)
4. Re-measure

Avoid guessing. The same code can behave very differently across GPU architectures.

Step 1 — Determine the Bottleneck (Theory)

Common bottlenecks:

• Memory-bound: high global memory traffic, low arithmetic intensity
• Compute-bound: lots of math per loaded byte, stalls less on memory
• Synchronization/latency-bound: frequent barriers or divergent branches
• Occupancy-limited: too many registers or shared memory per block

Key metrics (Nsight Compute concepts):

• Global load/store throughput and transactions
• Achieved occupancy / eligible warps per SM
• Stall reasons (e.g., Stall Not Selected, Stall Memory Dependency, etc.)
• Cache hit rates (L1/L2)

Code Example 1 — Kernel Fusion (Reduce Memory Traffic)

Two-pass pipelines often write intermediate results to global memory:

• Pass 1: y = f(x)
• Pass 2: out = g(y)

Fusion can keep intermediates in registers/shared memory and reduce DRAM traffic.

Fused kernel skeleton

#include <cuda_runtime.h>

__global__ void fused_f_g(const float* x, float* out, int N) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i >= N) return;

  // Intermediate stays in registers
  float y = x[i] * 2.0f;     // f(x)
  out[i]  = y + 1.0f;       // g(y)
}

Launch config

int threads = 256;
int blocks = (N + threads - 1) / threads;
fused_f_g<<<blocks, threads>>>(d_x, d_out, N);

Why this helps: fewer global writes/reads → higher effective bandwidth utilization.

Code Example 2 — Occupancy-Aware Launch Configuration

You can use CUDA occupancy APIs to choose block sizes that maximize “eligible warps” while respecting resource usage.

#include <cuda_runtime.h>
#include <cstdio>

__global__ void simple_kernel(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 << 20;

  int minGridSize = 0, blockSize = 0;
  size_t dynamicSMemSize = 0;

  // Use occupancy calculator to find a good block size.
  // (Note: for real work, consider using actual kernel attributes.)
  cudaOccupancyMaxPotentialBlockSize(
    &minGridSize,
    &blockSize,
    simple_kernel,
    dynamicSMemSize,
    0
  );

  int threads = blockSize;
  int blocks = (N + threads - 1) / threads;

  printf("Suggested block size: %d, grid size: %d\n", threads, blocks);
  // Launch: simple_kernel<<<blocks, threads>>>(...)
  return 0;
}

Practical Profiling Workflow (Theory)

1. Nsight Systems

   • Check CPU↔GPU overlap
   • Confirm kernels aren’t launched with tiny work sizes
   • Look for stalls between kernels (sync calls, host-GPU gaps)

2. Nsight Compute

   • Profile the main kernels
   • Inspect stall reasons
   • Compare variants: before/after tiling, fusion, block size changes

3. Validate

   • Regression test correctness
   • Benchmark multiple runs (warm-up, stable clocks)

Common Gotchas

• Profiling the wrong kernel: only optimize what dominates runtime.
• Ignoring launch overhead: too many tiny kernels can hurt more than you think.
• Over-optimizing for occupancy: sometimes fewer blocks but higher performance per block wins (e.g., better locality).
• Divergence in branches: warp divergence can inflate instruction count and reduce throughput.

Quick Checklist

• Profile with Nsight Systems to see “where time goes”
• Profile with Nsight Compute to see “why it stalls”
• Reduce DRAM traffic (fusion/tiling/vectorized loads)
• Tune block size with occupancy in mind
• Re-measure after each change