From CUDA Errors to Clean Builds: A Debugging Checklist

After years of working with CUDA at NVIDIA, I've compiled a practical checklist for debugging CUDA build issues. These are the hard-won lessons from maintaining production CUDA code in the Maxine SDK.

The Most Common Culprits

1. Missing Symbols and the -Wl,--whole-archive Trap

One of the most frustrating CUDA linking errors happens when device code gets stripped during linking. If you're seeing "undefined reference to __device__ function" errors, check if you're using --whole-archive incorrectly:

# Wrong - can cause device code stripping
target_link_libraries(myapp 
    -Wl,--whole-archive libcuda_kernels.a -Wl,--no-whole-archive)

# Right - use CMake's object libraries or proper CUDA linking
add_library(cuda_kernels OBJECT kernels.cu)
target_link_libraries(myapp $<TARGET_OBJECTS:cuda_kernels>)

2. The cudaGetLastError() Gotcha

CUDA errors are sticky - once an error occurs, it persists until explicitly cleared. This leads to mysterious failures where an error from one kernel shows up much later:

// Bad pattern - error might be from previous kernel
kernel1<<<grid, block>>>();
// ... lots of CPU code ...
kernel2<<<grid, block>>>();
checkCudaError(cudaGetLastError()); // Might report kernel1's error!

// Good pattern - check immediately after each kernel
kernel1<<<grid, block>>>();
checkCudaError(cudaGetLastError());
checkCudaError(cudaDeviceSynchronize());

kernel2<<<grid, block>>>();
checkCudaError(cudaGetLastError());

Build System Issues

3. NVCC Compilation Flags

Getting the right NVCC flags is crucial for performance and compatibility. Here's my go-to set:

set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} \
    -gencode arch=compute_70,code=sm_70 \
    -gencode arch=compute_75,code=sm_75 \
    -gencode arch=compute_80,code=sm_80 \
    -gencode arch=compute_86,code=sm_86 \
    -gencode arch=compute_86,code=compute_86 \
    --use_fast_math \
    -lineinfo")

Key points:

• Always include PTX (compute_XX) for forward compatibility
• --use_fast_math can give 10-15% speedup but changes semantics
• -lineinfo is essential for profiling with minimal overhead

4. Mixed C++/CUDA Compilation

When mixing C++ and CUDA code, symbol visibility can be tricky:

// In header file used by both C++ and CUDA
#ifdef __CUDACC__
    #define CUDA_CALLABLE __host__ __device__
#else
    #define CUDA_CALLABLE
#endif

class MyClass {
    CUDA_CALLABLE void process();
};

Runtime Debugging Techniques

5. Memory Access Violations

Use cuda-memcheck (or Compute Sanitizer for newer toolkits):

# Old toolkit
cuda-memcheck --leak-check full ./myapp

# CUDA 11.0+
compute-sanitizer --tool memcheck ./myapp

6. Kernel Launch Failures

Always validate your launch configuration:

void launchKernel(int dataSize) {
    int blockSize = 256;
    int gridSize = (dataSize + blockSize - 1) / blockSize;
    
    // Check limits
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, 0);
    
    if (gridSize > prop.maxGridSize[0]) {
        // Handle error - grid too large
    }
    
    if (blockSize > prop.maxThreadsPerBlock) {
        // Handle error - block too large
    }
    
    myKernel<<<gridSize, blockSize>>>(data, dataSize);
}

Performance Debugging

7. Warp Divergence

Profile with Nsight Compute to identify divergent branches:

ncu --metrics smsp__sass_average_branch_targets_threads_uniform.pct ./myapp

8. Memory Coalescing

Uncoalesced memory access can kill performance. Check with:

ncu --metrics l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,\
l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum ./myapp

The Complete Debugging Checklist

When facing a CUDA issue, go through this list systematically:

1. ✓ Check error after EVERY CUDA API call
2. ✓ Verify kernel launch configuration
3. ✓ Run with compute-sanitizer
4. ✓ Check for integer overflow in index calculations
5. ✓ Verify shared memory size doesn't exceed limits
6. ✓ Ensure proper stream synchronization
7. ✓ Check for race conditions (use __syncthreads())
8. ✓ Verify texture/surface binding (if using)
9. ✓ Profile with Nsight Compute for performance issues
10. ✓ Test with CUDA_LAUNCH_BLOCKING=1 for better error localization

Conclusion

CUDA debugging can be challenging, but with systematic approaches and the right tools, most issues become tractable. The key is to be methodical - check errors early and often, understand your hardware limits, and leverage the excellent profiling tools NVIDIA provides.

Remember: the compiler and runtime are trying to help you. Those cryptic error messages usually contain the exact information you need - you just need to know how to interpret them.