cuda-debugging

CUDA のデバッグ

目的

cuda-gdb による GPU スレッドの対話型検査、NVIDIA Compute Sanitizer によるメモリーエラーと競合の自動検出、GPU コアダンプの解析、デバイス側の printf、および一般的な CUDA ランタイムエラーコードのトリアージを通じて、エージェントが CUDA プログラムをデバッグできるようにします。

使用する場面

• cudaErrorIllegalAddress（700）またはデバイスでのセグメンテーション違反
• マルチスレッド GPU コードで断続的に発生する正確性の問題
• ワープ間、またはホストとデバイス間の競合のデバッグ
• cuda-gdb でカーネルコードを 1 行ずつ実行
• 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

cuda-gdb でソースレベルのステップ実行を行うには -G が必要です。サニタイザーは最適化ビルドでも動作しますが、-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. デバイス側の 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;
    }
}

# Buffer size for printf (default may truncate)
cuda-gdb) set cuda printf_buffer_size 16777216

出力を確認する前に cudaDeviceSynchronize() でフラッシュしてください。すべてのスレッドから大量の printf を行うと、バッファーがあふれます。

5. エラーコードのトリアージ

コード	名前	一般的な原因
700	cudaErrorIllegalAddress	範囲外アクセス、解放後の使用、不正なポインター
701	cudaErrorLaunchOutOfResources	ブロックあたりの共有メモリーまたはレジスターが多すぎる
702	cudaErrorLaunchTimeout	無限ループ、TDR ウォッチドッグ（Windows／Linux のデフォルト）
719	cudaErrorLaunchFailure	カーネル内のアサート、スタックオーバーフロー

// Always check after launch
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 コアダンプ

# Enable coredump (driver 450+)
export CUDA_ENABLE_COREDUMP_ON_EXCEPTION=1
export CUDA_COREDUMP_FILE=/tmp/cuda_coredump_%h.%p

./app_san   # crash generates dump

# Analyze with 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. デバッグの判断ツリー

Crash or wrong results?
├── Consistent wrong values → logic bug; use printf or cuda-gdb
├── Intermittent / depends on size → OOB or race
│   ├── compute-sanitizer --tool memcheck
│   └── compute-sanitizer --tool racecheck
├── Hang / timeout 702 → infinite loop or barrier mismatch
│   └── synccheck; audit __syncthreads paths
└── Works in debug (-G), fails in release → uninitialized mem or race
    └── initcheck + racecheck on release build

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 がアトミック操作で誤検知する	アトミックでない RMW	atomicAdd／atomicCAS を使用する
アタッチに失敗する	プロセスが CUDA コンテキスト内にない	最初の cudaMalloc の後でブレークする
Windows で TDR タイムアウトが発生する	カーネルの実行時間が長い	カーネルを分割するか、cudaDeviceSetLimit または regedit で TDR を設定する

関連スキル

• skills/gpu/cuda — カーネルのパターン、メモリー階層、起動設定
• skills/gpu/cuda-profiling — 正確性を確認した後のパフォーマンス
• skills/gpu/gpu-memory-model — 競合とコアレッシングの理解
• skills/debuggers/gdb — cuda-gdb と共通するホスト側の GDB コマンド
• skills/runtimes/sanitizers — ホストコード向けの ASan／TSan の概念
• skills/kernel/kernel-debugging — ドライバーレベルの問題向け kgdb／kprobes