Debugging & Correctness

Compile with -lineinfo to get source line numbers.

33.1 compute-sanitizer

Compile with -lineinfo to get source line numbers.

compute-sanitizer ./app                        # memcheck: out-of-bounds, misaligned, leaks (with --leak-check full)
compute-sanitizer --tool racecheck ./app       # shared-memory data races
compute-sanitizer --tool initcheck ./app       # uninitialized global memory reads
compute-sanitizer --tool synccheck ./app       # invalid synchronization (e.g., divergent __syncthreads)

33.2 cuda-gdb

Nsight Visual Studio Edition and Nsight VS Code Edition provide GUI debugging.

nvcc -g -G app.cu -o app
cuda-gdb ./app
(cuda-gdb) break myKernel
(cuda-gdb) run
(cuda-gdb) info cuda threads
(cuda-gdb) cuda block 1 thread 5            # switch focus to a specific thread
(cuda-gdb) print data[threadIdx.x]
(cuda-gdb) next / step / continue

33.3 Debugging strategies

Check every API call and kernel launch (CUDA_CHECK).

Set CUDA_LAUNCH_BLOCKING=1 to make launches synchronous (pinpoints the failing kernel).

Use printf in kernels (limit output; buffer size via cudaDeviceSetLimit(cudaLimitPrintfFifoSize, ...)).

assert() in device code (#include <cassert>).

Compare against a CPU reference implementation with tolerances (floating-point order differs!).

Test edge cases: n = 0, 1, non-multiples of block size, very large sizes.

Watch for: missing __syncthreads(), races, integer overflow in index math (use size_t/int64_t for large arrays), uninitialized shared memory, reading results before synchronization.

33.4 Floating-point determinism

Parallel reductions change summation order → small result differences vs the CPU; atomics make results non-deterministic between runs. Use deterministic algorithms (fixed-order trees), Kahan summation, or higher-precision accumulation when reproducibility matters.