Two tools address the correctness of CUDA programs:
cuda-gdb, a source-level debugger for device as well as
host code, and compute-sanitizer, a suite of runtime
checkers for the memory and synchronization errors that are easy to
introduce in parallel code and hard to find by inspection.
cuda-gdb extends the GNU debugger gdb with
the ability to set breakpoints in __global__ and
__device__ code, single-step a warp, and inspect device
memory and variables. Debugging device code requires that the
application be built with device debug information, which
nvcc emits when passed -G (the companion to
the host-side -g); because -G disables most
optimizations, debug builds run more slowly and can hide or expose bugs
that depend on timing. A session resembles an ordinary gdb
session – break, run, continue,
next, print – with extra commands for
navigating the thread hierarchy:
cuda thread (0,0,0) block (1,0,0) moves the focus to a
specific thread, and info cuda threads,
info cuda warps, and info cuda kernels report
the state of the running grid. On Pascal and later GPUs,
cuda-gdb can debug a kernel on the same GPU that drives the
display, so a dedicated second GPU is no longer required.
compute-sanitizer, which replaced the older
cuda-memcheck, runs a program under one of four checkers,
selected with --tool:
memcheck, the default, reports out-of-bounds and
misaligned accesses to global, local, and shared memory, and, with
--leak-check full, allocations that were never
freed.
racecheck detects shared memory data races –
accesses from different threads to the same location that are not
separated by a barrier.
initcheck flags reads of global memory that was
never written.
synccheck catches invalid use of
__syncthreads() and related barriers, such as a barrier
that not all threads in the block can reach.
A typical invocation is
compute-sanitizer --tool racecheck ./myApp. Each report
names the kind of error, the offending kernel, and the block and thread
that triggered it; building with -lineinfo adds the source
file and line, which turns most reports into a direct pointer at the
bug. compute-sanitizer needs no special build and runs
memcheck at close to full speed, though
racecheck and initcheck are considerably
slower.
Not every bug needs a debugger. Two facilities compiled into device
code let a kernel report on itself from the inside:
printf() and assert(), both available since
the Fermi architecture (SM 2.0).
printf() is callable from __global__ and
__device__ code and accepts the same format strings and
conversions as the host function. What differs is the path the output
takes to the console. Each call copies its format string and arguments
into a circular buffer in device global memory; that buffer is drained
to the host’s stdout only at defined synchronization points
– when the kernel completes, at cudaDeviceSynchronize(),
and at blocking driver calls – never while the kernel is still running.
A kernel that hangs or faults before reaching such a point produces no
output at all, so a missing line is itself a clue. The buffer has a
fixed size, one megabyte by default, set with
cudaDeviceSetLimit(cudaLimitPrintfFifoSize, bytes); when it
fills, the oldest contents are overwritten, so a grid that prints
without restraint can lose the very lines it needs. Restraint matters
for a second reason: every thread that reaches the call executes it, so
an unguarded printf() in a large grid emits one line per
thread. Useful output almost always guards the call on a chosen thread –
if (blockIdx.x == 0 && threadIdx.x == 0) – or on
the few threads whose state is in question. Output from different
threads and blocks interleaves in an unspecified order; only a single
thread’s call is emitted as a unit.
assert() evaluates a condition in device code and, when
it is false, prints the expression, source file, and line, halts the
kernel, and puts the context into an error state – every subsequent CUDA
call returns cudaErrorAssert until the context is
destroyed. It pins a violated invariant to the exact thread and line
that broke it, and, like the host assert(), it compiles to
nothing under -DNDEBUG.