cuda-debugging

CUDA Debugging

Purpose

Guide agents through debugging CUDA programs with cuda-gdb for interactive GPU thread inspection, NVIDIA Compute Sanitizer for automated memory and race detection, GPU core dump analysis, device-side printf, and triaging common CUDA runtime error codes.

When to Use

• cudaErrorIllegalAddress (700) or segmentation fault on device
• Intermittent correctness failures in multi-threaded GPU code
• Debugging race conditions between warps or between host and device
• Stepping through kernel code line-by-line with cuda-gdb
• Validating uninitialized memory reads with initcheck
• Kernel hang or cudaErrorLaunchTimeout (702)

Workflow

1. Build for debugging

# Debug build — disables optimizations, enables device debug
nvcc -G -g -O0 -arch=sm_80 -o app_debug main.cu

# Sanitizer-friendly build (lineinfo helps reports)
nvcc -lineinfo -g -O2 -arch=sm_80 -o app_san main.cu

-G is required for cuda-gdb source-level stepping. Sanitizers work with optimized builds but -G gives clearer line numbers.

2. Compute Sanitizer — automated checks

# Memory errors (OOB, misaligned, leak)
compute-sanitizer --tool memcheck ./app_san

# Shared memory and global memory races
compute-sanitizer --tool racecheck ./app_san

# Uninitialized memory reads
compute-sanitizer --tool initcheck ./app_san

# Synchronization errors (missing __syncthreads)
compute-sanitizer --tool synccheck ./app_san

# Verbose with source correlation
compute-sanitizer --tool memcheck --show-reachable=yes --log-file san.log ./app_san

Typical memcheck output:

======== Invalid __global__ write of size 4
========     at 0x1a0 in vector_add(vector_add.cu:12)
========     by thread (0,0,0) in block (0,0,0)
======== Address 0x7f... is out of bounds

3. cuda-gdb interactive session

# Launch under cuda-gdb
cuda-gdb ./app_debug

# Or attach to running process
cuda-gdb -p <pid>

Essential commands:

# Break at kernel entry
(cuda-gdb) break vector_add
(cuda-gdb) run

# Focus on GPU threads
(cuda-gdb) info cuda kernels
(cuda-gdb) cuda kernel 0
(cuda-gdb) cuda thread (0,0,0)   # block (x,y,z), thread (x,y,z)

# Inspect device memory
(cuda-gdb) print data[i]
(cuda-gdb) x/10f d_ptr

# Step in kernel
(cuda-gdb) cuda step
(cuda-gdb) cuda next

# All threads in block
(cuda-gdb) info cuda threads
(cuda-gdb) cuda thread (0,0,5)