Profiling for Success: Maximizing Throughput in CUDA Applications#

2.1 Kernel Definition#

In CUDA, a kernel is the function you launch on the GPU. Each kernel runs in a multi-threaded environment that you define using a grid, which is subdivided into thread blocks. For example:

1
__global__ void vectorAdd(const float* A, const float* B, float* C, int N) {
2
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
3
    if (idx < N) {
4
        C[idx] = A[idx] + B[idx];
5
    }
6
}

In this example, each thread adds one element from array A to one element from array B. When you launch the kernel, you specify the dimensions of the grid and blocks, for instance:

1
int blockSize = 256;
2
int gridSize = (N + blockSize - 1) / blockSize;
3
vectorAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);

2.2 Thread Hierarchy and Memory#

Each thread block has blockDim.x * blockDim.y * blockDim.z threads, grouped into warps of 32 threads each. Performance is heavily dependent on how these threads access memory (coalescing) and how many active threads per SM (occupancy) are possible.

3.1 Nsight Systems#

Nsight Systems offers a timeline of CPU and GPU activity. You can detect synchronization points, concurrent kernel execution, or pipeline stalls. This helps you see if your application is GPU-bound or CPU-bound.

3.2 Nsight Compute#

Nsight Compute drills down on individual kernels, displaying details like memory throughput, instruction throughput, stalls, and multiprocessor occupancy. It also suggests specific optimization strategies, such as improving memory coalescing or reducing shared memory bank conflicts.

3.3 Command-Line Profiling#

Command-line options for Nsight Systems and Nsight Compute let you script the profiling process. For example, you might run:

1
nsys profile -o app_timeline ./my_gpu_app

or

1
ncu -o kernel_analysis ./my_gpu_app

Both commands generate profiling results you can analyze in GUI tools or through command-line summaries.

4.1 Memory Access Patterns#

A common source of inefficiency is suboptimal memory access. GPUs handle large memory transactions more effectively if threads in a warp address consecutive memory locations. When accesses are scattered, memory bandwidth is underutilized.

Tips for optimal memory access:

1. Ensure that each thread in a warp accesses contiguous memory.
2. Use shared memory to optimize repeated accesses to the same data.
3. Minimize uncoalesced writes, which can severely degrade performance.

4.2 Kernel Launch Configuration#

Choosing the right number of thread blocks and block size can significantly influence occupancy and performance. More threads can keep the GPU busy, but too many can also result in more register usage per thread, reducing occupancy.

You can start with a block size of 128 or 256 threads and adjust according to performance. Tools like Nsight Compute will show you the achieved occupancy. Aim for around 60-100% SM occupancy (depending on the kernel’s resource usage).

4.3 Loop Unrolling and Compiler Optimizations#

The NVIDIA compiler (nvcc) supports various optimizations:

• Loop unrolling: Can reduce overhead and increase pipeline efficiency.
• Inline expansion: Reduces function call overhead but can increase register usage.
• Constant propagation: Replaces known compile-time constants.

Choose appropriate compiler flags, such as -O3 for maximum optimization, but beware of increased compilation time and potential register usage.

5.1 Identifying Bottlenecks#

When analyzing a kernel in Nsight Compute, you might notice that memory operations are hitting near 100% utilization but the SM is underutilized. This suggests you’re memory-bound, and focusing on global memory access patterns or caching might help. Conversely, if memory utilization is low but your kernel has high instruction throughput, you may be compute-bound and should consider algorithmic optimizations or specialized instructions (e.g., using Tensor Cores if your operations warrant it).

5.2 Annotating Your Code for Better Profiling#

Annotated source code can aid in pinpointing specific regions of interest. Using the NVTX (NVIDIA Tools Extension) library, you can create ranges in your code:

1
#include <nvToolsExt.h>
2

3
void myFunction() {
4
    nvtxRangePush("MyFunction");
5
    // ... code you want to profile ...
6
    nvtxRangePop();
7
}

Nsight Systems or Nsight Compute can display these annotations in the timeline, making it easier to correlate code sections with performance events.

6.1 Global Memory Coalescing#

As mentioned, coalesced accesses occur when threads in the same warp access consecutive addresses. This results in a single or minimal number of memory transactions rather than many. Below is a table summarizing good vs. poor memory access patterns:

6.2 Shared Memory: Tuning and Bank Conflicts#

Shared memory can be thought of as a user-managed cache. Efficient use can dramatically speed up operations that require repeated access to the same data subset. However, shared memory is organized in banks, and if two or more threads in a warp access the same bank, performance degrades due to bank conflicts.

How to reduce bank conflicts:

• Align data structures in shared memory to warp boundaries when possible.
• Use padding to avoid multiple threads mapping to the same bank.

6.3 Constant and Texture Memory#

For read-only data that doesn’t change throughout the kernel’s execution, constant and texture memory can be advantageous. They offer caching benefits and can handle broadcast reads efficiently if many threads read the same location. Profile the frequency of these reads to see if the overhead of using specialized memory spaces is worthwhile.

7.1 Calculating Occupancy#

CUDA provides an API for calculating theoretical occupancy:

1
cudaError_t status;
2
int blockSize = 256;
3
int minGridSize, blockSizeChosen;
4
status = cudaOccupancyMaxPotentialBlockSize(
5
    &minGridSize,
6
    &blockSizeChosen,
7
    vectorAdd,
8
    0,
9
    0
10
);

You can compare the chosen block size with your chosen block size, or adjust resource usage (e.g., shared memory per block) to affect occupancy. The Nsight Compute Occupancy analysis also reveals if you’re limited by threads, registers, or shared memory.

7.2 Balancing Register Usage#

Register usage is critical to GPU performance. Each SM has a limited register file; excessive usage reduces the number of concurrent warps. While more registers per thread can improve performance by reducing global memory accesses, going too far can degrade occupancy.

Tip: Use -maxrregcount to limit registers, or inspect your kernel’s register usage in Nsight Compute. Sometimes, adjusting your code to reduce variable usage can reclaim registers for more threads.

8.1 Warp Divergence#

Because threads in a warp execute in lockstep, if they take different branches, the warp must serialize those branches. This is known as warp divergence and can destroy parallel efficiency. For instance, code like this can cause divergence:

1
if (threadIdx.x % 2 == 0) {
2
    // Do even work
3
} else {
4
    // Do odd work
5
}

When half the warp does one set of instructions and half does another, the final throughput suffers. If possible, restructure your algorithm to minimize conditional branching within warps.

8.2 Predication vs. Branching#

The compiler often uses predication, which executes both paths of a branch but discards results for threads that shouldn’t execute a path. This eliminates branch overhead but still incurs the cost of running instructions. Profile your kernel to see if predication is more efficient than actual branching in your scenario.

9.1 Baseline Profile#

Profiling this naive approach will often reveal:

• High global memory traffic.
• Lower-than-ideal occupancy if N is large and loop unrolled in registers.
• A massive number of redundant accesses to the same elements.

9.2 Shared Memory Optimization#

A classic approach to optimizing matrix multiplication is to load tiles of A and B into shared memory. Each block computes a tile of the resultant C matrix. An example of a shared memory tiles approach:

1
#define TILE_WIDTH 16
2

3
__global__ void matMulShared(const float* A, const float* B, float* C, int N) {
4
    __shared__ float As[TILE_WIDTH][TILE_WIDTH];
5
    __shared__ float Bs[TILE_WIDTH][TILE_WIDTH];
6

7
    int tx = threadIdx.x, ty = threadIdx.y;
8
    int row = blockIdx.y * TILE_WIDTH + ty;
9
    int col = blockIdx.x * TILE_WIDTH + tx;
10

11
    float sum = 0.0f;
12
    for (int phase = 0; phase < (N / TILE_WIDTH); phase++) {
13
        As[ty][tx] = A[row * N + (phase * TILE_WIDTH + tx)];
14
        Bs[ty][tx] = B[(phase * TILE_WIDTH + ty) * N + col];
15
        __syncthreads();
16

17
        for (int k = 0; k < TILE_WIDTH; k++) {
18
            sum += As[ty][k] * Bs[k][tx];
19
        }
20
        __syncthreads();
21
    }
22
    if (row < N && col < N) {
23
        C[row * N + col] = sum;
24
    }
25
}

Key improvements:

• Each element of A and B is loaded from global memory only once per tile.
• The inner dot product loop is performed in shared memory, reducing global memory bandwidth demand.

9.3 Profiling the Optimized Kernel#

Using Nsight Compute, you’ll see:

• Increased data reuse: The shared memory hits go up, global memory usage goes down.
• Reduced memory bandwidth usage for large N.
• Possible higher occupancy if tile sizes match hardware constraints.

If performance is still limited, consider further improvements like using half-precision (FP16) or Tensor Cores (on suitable GPus) for matrix multiplication, which can drastically increase throughput.

10.1 Asynchronous Streams#

CUDA streams allow you to launch kernels asynchronously, potentially overlapping computation and data transfers. By default, all operations run in the default stream, which is sequential. Creating multiple streams can improve concurrency:

1
cudaStream_t stream1, stream2;
2
cudaStreamCreate(&stream1);
3
cudaStreamCreate(&stream2);
4

5
kernelA<<<blocksA, threadsA, 0, stream1>>>(...);
6
kernelB<<<blocksB, threadsB, 0, stream2>>>(...);

If kernelA and kernelB are independent, they can run concurrently on different SM resources, leading to better GPU utilization. Nsight Systems shows stream timelines, enabling you to see if kernels are truly overlapping.

10.2 Data Transfers#

Memory transfers between the host and GPU are another common bottleneck. Use pinned (page-locked) memory for faster transfers, and schedule them in parallel with kernel execution (via separate streams). Profiling will reveal if your GPU is idle waiting for data transfers.

12.1 Mixed Precision#

On GPUs with Tensor Cores, mixed-precision arithmetic (FP16 or TF32 on Ampere) can achieve higher throughput than standard FP32 or FP64 operations. This is critical in domains like deep learning, where slight rounding differences are often acceptable. Profiling can confirm if your operations are bound by numeric precision or if they can leverage Tensor Cores fully.

12.2 Using Intrinsics#

CUDA provides intrinsics like __fmaf_rn(a, b, c) for multiply-add operations. These can sometimes outperform sequence-of-operations equivalents. Profiling will tell you if calls to such intrinsics reduce instruction latency and improve throughput. Use them judiciously to match the precision and rounding mode you need.

13.1 Identifying Bottlenecks#

A timeline in Nsight Systems might reveal that each convolution call dominates the application’s runtime. Nsight Compute might show:

• High pipeline stalls due to repeated access to the same input elements.
• High global memory transactions as the kernel reads from input for each thread.

13.2 Optimization Approaches#

1. Use shared memory to load a tile of the input image.
2. Convert repeated boundary checks (min, max) into masked operations or pad the image beforehand to avoid conditionals inside the loop.
3. For small kernels, consider warp-level intrinsics to handle partial sums.
4. Use texture fetches if 2D spatial locality is strong.

Profiling these optimizations iteratively helps you determine which approach yields the best speedup. In many image-based kernels, shared memory tiling can lead to significant performance gains.