cuda-debugging

用途

本技能用于指导代理调试 CUDA 程序，涵盖使用 cuda-gdb 进行交互式 GPU 线程检查、使用 NVIDIA Compute Sanitizer 自动检测内存与竞态问题、GPU core dump 分析、device 端 printf 输出，以及排查常见 CUDA 运行时错误码。

适用场景

cudaErrorIllegalAddress (700) 或设备端段错误
多线程 GPU 代码中间歇性正确性失败
调试 warp 之间或主机与设备之间的竞态条件
使用 cuda-gdb 逐行单步调试内核代码
使用 initcheck 校验未初始化内存读取
内核挂起或 cudaErrorLaunchTimeout (702)

工作流程

1. 构建调试版本

调试构建可禁用优化并启用设备调试：

# 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 是 cuda-gdb 源码级单步调试所必需的。Sanitizer 也可用于优化构建，但 -G 能提供更清晰的行号。

2. Compute Sanitizer — 自动化检查

# 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

典型的 memcheck 输出：

======== 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 交互式会话

启动方式：

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

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

常用命令：

# 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)

4. Device printf

__global__ void debug_kernel(float *data, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        if (i < 5)  // limit output
            printf("thread %d: data[%d] = %f\n", i, i, data[i]);
        data[i] *= 2.0f;
    }
}

设置 printf 缓冲区大小（默认可能截断）：

(cuda-gdb) set cuda printf_buffer_size 16777216

检查输出前需用 cudaDeviceSynchronize() 刷新。所有线程的大量 printf 会淹没缓冲区。

5. 错误码排查

代码	名称	常见原因
700	cudaErrorIllegalAddress	越界访问、释放后使用、错误指针
701	cudaErrorLaunchOutOfResources	每块共享内存或寄存器过多
702	cudaErrorLaunchTimeout	无限循环、TDR 看门狗（Windows/默认 Linux）
719	cudaErrorLaunchFailure	内核断言、栈溢出

启动后始终检查错误：

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));

6. GPU core dump

启用 coredump（驱动 450+）：

export CUDA_ENABLE_COREDUMP_ON_EXCEPTION=1
export CUDA_COREDUMP_FILE=/tmp/cuda_coredump_%h.%p

./app_san   # crash generates dump

使用 cuda-gdb 分析：

cuda-gdb ./app_san /tmp/cuda_coredump_hostname.pid
(cuda-gdb) cuda coredump load /tmp/cuda_coredump_hostname.pid
(cuda-gdb) bt
(cuda-gdb) info cuda kernels

7. 调试决策树

崩溃还是结果错误？

一致的错误值 → 逻辑 bug；使用 printf 或 cuda-gdb
间歇性 / 依赖规模 → 越界或竞态
- compute-sanitizer --tool memcheck
- compute-sanitizer --tool racecheck
挂起 / 超时 702 → 无限循环或屏障不匹配
- synccheck；审查 __syncthreads 路径
调试版（-G）正常，发布版失败 → 未初始化内存或竞态
- 在发布版上运行 initcheck + racecheck

8. 多 GPU 与 MIG 说明

# Isolate GPU
CUDA_VISIBLE_DEVICES=0 compute-sanitizer --tool memcheck ./app

# MIG instances appear as separate devices
nvidia-smi -L

常见问题

现象	原因	修复
cuda-gdb 无法在内核断点	未用 -G 构建	用 nvcc -G -g -O0 重新构建
Sanitizer 无报错但崩溃依旧	异步错误延迟	在内核后添加 cudaDeviceSynchronize()
printf 无输出	缓冲区满或未同步	限制打印；增大缓冲区；同步
racecheck 对原子操作误报	非原子读改写	使用 atomicAdd/atomicCAS
attach 失败	进程不在 CUDA 上下文中	在首次 cudaMalloc 后断点
Windows 下 TDR 超时	内核运行过长	拆分内核；cudaDeviceSetLimit 或修改注册表 TDR

技能介绍

用途