Cuda Debugging

CUDA debugging

Contract

Inputs

1. Symptom (required): crash, wrong values, hang, or timeout, and when it appears.
2. Kernel source (required): the .cu files under suspicion.
3. Toolkit (required): nvcc --version and compute-sanitizer --version. Grounded current stable: CUDA Toolkit 13.3 Update 1.
4. Target GPU and any MIG layout (optional): nvidia-smi -L.

Procedure

1. Build for the tool. Source-level stepping in cuda-gdb needs nvcc -G -g -O0. Sanitizers accept optimized builds and report clearer lines with -lineinfo -g -O2:

nvcc -G -g -O0 -arch=sm_80 -o app_debug main.cu
nvcc -lineinfo -g -O2 -arch=sm_80 -o app_san main.cu

Done when: both binaries exist and -G is absent from the sanitizer build.

2. Run the Compute Sanitizer tools in order of suspicion. Each tool slows execution sharply; run one at a time:

compute-sanitizer --tool memcheck ./app_san    # out-of-bounds, misaligned, leaks
compute-sanitizer --tool racecheck ./app_san   # shared and global memory races
compute-sanitizer --tool initcheck ./app_san   # uninitialized reads
compute-sanitizer --tool synccheck ./app_san   # missing or misaligned __syncthreads
compute-sanitizer --tool memcheck --log-file san.log ./app_san

A typical memcheck hit names the kernel, line, thread, and address:

======== 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

Done when: the failing tool's report is captured or all four tools pass.

3. Step through the kernel in cuda-gdb when static reports are not enough:

cuda-gdb ./app_debug

(cuda-gdb) break vector_add
(cuda-gdb) run
(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)
(cuda-gdb) print data[i]
(cuda-gdb) next                  # standard step commands work in kernels
(cuda-gdb) info cuda threads

Done when: the suspect thread's state at the fault is inspected.

4. Use device printf for cheap tracing inside kernels. Cap the output to a few threads, and flush with cudaDeviceSynchronize() before reading. The printf buffer is finite; set cuda printf_buffer_size in cuda-gdb raises it when needed:

if (i < n && i < 5)
    printf("thread %d: data[%d] = %f\n", i, i, data[i]);

Done when: printed values confirm or refute the hypothesis.

5. Check every launch and sync for errors before diagnosing deeper:

kernel<<<grid, block>>>(args);
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess)
    fprintf(stderr, "launch: %s\n", cudaGetErrorString(err));
cudaDeviceSynchronize();
err = cudaGetLastError();
if (err != cudaSuccess)
    fprintf(stderr, "exec: %s\n", cudaGetErrorString(err));

Triage the returned code:

Done when: the code is mapped to a cause and the corresponding tool from step 2 confirms it.

6. Capture a core dump when the crash cannot be reproduced under a debugger. Export the two variables, run until the fault, then load the dump:

export CUDA_ENABLE_COREDUMP_ON_EXCEPTION=1
export CUDA_COREDUMP_FILE=/tmp/cuda_coredump_%h.%p
./app_san
cuda-gdb ./app_san /tmp/cuda_coredump_hostname.pid

Done when: bt and info cuda kernels on the dump show the faulting wave.

7. Follow the decision path. Consistent wrong values point to a logic bug; use printf or cuda-gdb. Size-dependent or intermittent failures point to out-of-bounds or races; run memcheck and racecheck. A hang points to an infinite loop or a barrier mismatch; run synccheck and audit every __syncthreads path. A program that works under -G and fails optimized points to uninitialized memory or a race; run initcheck and racecheck on the release build. An async error that appears late needs a cudaDeviceSynchronize() right after the suspect launch. Done when: the observed failure class matches exactly one branch and its evidence exists.
8. For multi-GPU or MIG hosts, isolate the device under test with CUDA_VISIBLE_DEVICES=0 (a MIG UUID for MIG instances) and confirm the visible set with nvidia-smi -L. Done when: the sanitizer or debugger runs against one named device.

Failure and recovery

Output

A triage verdict: the failure class, the tool and evidence that named it (report excerpt, thread state, or error code), the defect location, and the recommended fix. Sanitizer logs and dumps stay in the project tree or /tmp for the caller to delete.