【入門編】GPUプログラミングの壁を越える：CUDA/ROCm環境でGDBを利用したカーネルコードのデバッグ手法 – デバッグ・コード品質・テストツール生産性向上バイブル

こんにちは！日々のハードなデバッグ作業、本当にお疲れ様です。
CPUのコードなら `printf` や使い慣れたIDEのブレークポイントでサクッと解決できるのに、いざ「GPU（CUDA / ROCm）」の世界に足を踏み入れた途端、何が起きているのかさっぱり分からなくなって途方に暮れた経験はありませんか？

「カーネル内でセグメンテーション違反（GPU側では `cudaErrorIllegalAddress` など）が起きたのに、どのスレッドで落ちたのか分からない」
「変数の値を見ようとしても、数千・数万の並列スレッドが同時に動いているから頭がクラクラする」

GPUプログラミングを始めたエンジニアの多くが、この「CPUとGPUの壁」の高さに絶望します。しかし、安心してください。今日ここでマスターする GDB（NVIDIAなら `cuda-gdb`、AMD ROCmなら `roc-gdb`） の実践的なテクニックを身につければ、その壁は驚くほど低くなります。

これをマスターすれば、毎日のコーディングが劇的に楽になりますよ。さあ、一緒にGPUの内部を覗き込む冒険に出かけましょう！

—

1. なぜGPUのデバッグは難しいのか？（アーキテクトからの前提知識）

通常のCPUプログラムは、せいぜい数十スレッド（マルチコア）がOSによってスケジュールされながら動いています。そのため、ブレークポイントで止めれば「どのスレッドがどこにいるか」を直感的に把握できます。

しかし、GPUは違います。1つのカーネル起動（Kernel Launch）に対し、数万から数百万の「スレッド」が、単一命令複数データ（SIMT）アーキテクトのもとで超並列実行されます。
つまり、デバッガから見ると「数万個のCPUコアが全く同じコードを異なるデータで一斉に走らせている状態」です。このカオスを制御するために、低レイヤデバッガの内部アーキテクトたちは、「フォーカス（Focus）」という概念を用意しました。

GDBを使いこなすとは、この数万の群れ（グリッドとブロック）の中から「今、バグを引き起こしている張本人のスレッド」をピンポイントで捕獲する技術に他なりません。

—

2. 開発環境の基礎セットアップ：ただコンパイルするな

GPU用GDB（CUDA環境なら `cuda-gdb`）でコードをステップ実行するためには、コンパイラに対して「GPUの機械語（SASS）だけでなく、デバッグ用のメタデータ（DWARF格式）をバイナリに埋め込め」と明示的に指示する必要があります。

以下の極めてシンプルな「HelloWorld」的なCUDAコード（ベクトル加算）を例に、正しいコンパイル手順を見ていきましょう。

サンプルコード: `vector_add.cu`

include

// GPU上で実行されるカーネル関数
__global__ void vectorAdd(const float A, const float B, float C, int numElements) {
int i = blockDim.x blockIdx.x + threadIdx.x;
if (i < numElements) { // わざと複雑な演算や、デバッグで追いかけやすい処理を記述 C[i] = A[i] + B[i]; } } int main(void) { int numElements = 512; size_t size = numElements sizeof(float); // ホスト（CPU）側のメモリ確保と初期化 float h_A = (float )malloc(size); float h_B = (float )malloc(size); float h_C = (float )malloc(size); for (int i = 0; i < numElements; ++i) { h_A[i] = (float)i; h_B[i] = (float)(i 2); } // デバイス（GPU）側のメモリ確保 float d_A = NULL; cudaMalloc((void )&d_A, size); float d_B = NULL; cudaMalloc((void )&d_B, size); float d_C = NULL; cudaMalloc((void )&d_C, size); // データの転送 cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice); cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice); // カーネルの起動（1ブロック、512スレッド） int threadsPerBlock = 512; int blocksPerGrid = 1; vectorAdd<<>>(d_A, d_B, d_C, numElements);

// 結果の回収
cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

// 検証
printf(“Test PASSED\n”);

// クリーンアップ
free(h_A); free(h_B); free(h_C);
cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
return 0;
}

命運を分けるコンパイルフラグ

このコードをビルドする際、普通の `nvcc vector_add.cu -o vector_add` ではデバッグ情報が落ちてしまいます。以下のようにコンパイルしてください。

-g: ホスト（CPU）側のデバッグ情報を付与
-G: デバイス（GPU）側のデバッグ情報（Device Debugging）を有効化
-O0: コンパイラによる最適化を無効化（変数が消えるのを防ぐ）
nvcc -g -G -O0 vector_add.cu -o vector_add

> アーキテクトからのワンポイントアドバイス：
> `-G` フラグをつけると、GPU上のコード実行速度は劇的に低下します（ハードウェアの最適化機能が一部バイパスされるため）。しかし、バグを取るための初期段階では「必ず `-g -G -O0`」を鉄則として体に覚え込ませてください。

—

3. 実戦：`cuda-gdb` を起動し、カーネルをキャッチする

それでは、実際にデバッガをアタッチして動作確認を行います。ターミナルを開いてください。

cuda-gdbを起動し、対象のバイナリを読み込ませる
$ cuda-gdb ./vector_add

GDBのプロンプト（`(cuda-gdb)`）が立ち上がったら、まずはGPUのコード（カーネル関数）にブレークポイントを仕掛けます。CPUと同じように `break` 命令が使えます。

(cuda-gdb) break vectorAdd
Breakpoint 1 at 0x3d0: file vector_add.cu, line 6.

プログラムを実行（`run`）します。

(cuda-gdb) run
Starting program: /path/to/vector_add
[Thread debugging using libthread_db enabled]
Using host libthread_db library “/lib/x86_64-linux-gnu/libthread_db.so.1”.
[New Thread 0x7ffff7fba700 (LWP 12345)]
[Switching to CUDA Focus]

おっ！プログラムが動き出し、GPUのカーネル関数にヒットした瞬間に自動でフォーカスが切り替わりました。

—

4. 並列スレッドの群れを支配する：スレッドとフォーカスの操作

ここで `backtrace` やスレッド一覧を見てみましょう。

(cuda-gdb) info threads

実行すると、数百ものCUDAスレッド（Thread 1, Thread 2, … Thread 512）がずらりと表示されます。この中で、今自分がどのスレッドを見ているのかを確認・変更するのが `cuda thread` コマンドです。

現在フォーカス当たっているスレッドの情報を確認
(cuda-gdb) cuda thread
[Current CUDA Focus: Device 0, Block 0, Thread (42,0,0), PCI Bus ID 0000:01:00.0]

今、ブロック0の、X軸42番目のスレッドにフォーカスが当たっています。
例えば、「スレッド番号 100 の世界では、変数 `i` や `A[i]` がどうなっているか」を直接覗いてみましょう。

フォーカスを特定のスレッド（例: Block 0, Thread 100）に移動
(cuda-gdb) cuda thread (0,0,100)

[Switching focus to CUDA Thread 0,0,100, grid 1, block (0,0,0), thread (100,0,0), device 0, sm 0, warp 3, lane 4]

そのスレッドから見たローカル変数を表示
(cuda-gdb) print i
$1 = 100
(cuda-gdb) print A[i]
$2 = 100

このように、「数万のスレッドが並列実行されている空間から、任意の座標（Block / Thread）を指定して顕微鏡のように覗き込む」ことができるのが、GPUデバッグの最大の強みであり、快感です。

—

5. デバイスメモリを一網打尽にダンプするテクニック

GPUプログラミングで最も頭を悩ませるのが「ホストとデバイス間のメモリ不整合」や「GPU側で書き換えられたデータの破損」です。

通常の `print` では、デバイス側（GPUのVRAM上）にあるポインタの中身を直接CPU側のように人間が読みやすい形式で展開できない場合があります。そんなときは、GDBの `x`（Examine Memory）コマンドを使いこなします。

例えば、デバイス側メモリ `d_C` の先頭から10個の浮動小数点数（float）をダンプしたい場合：

デバイスメモリのポインタ d_C のアドレスを確認
(cuda-gdb) print d_C
$3 = (float ) 0x7f84c0000000

10個分のfloat (4バイト 10) を16進数/浮動小数点形式でメモリダンプ
(cuda-gdb) x/10f d_C
0x7f84c0000000: 0 1 2 3
0x7f84c0000010: 4 5 6 7
0x7f84c0000020: 8 9

一撃でGPU上のメモリが綺麗に可視化されました！
もしここで意図しないゴミデータが入っていれば、どのカーネルスレッドが書き損じたのかを `cuda thread` を切り替えながら追跡すれば、バグの原因は秒速で特定できます。

—

6. まとめ：GPUデバッグは「怖くない」

ここまで、CUDA/ROCm環境におけるGDB（`cuda-gdb`）の基本的なアプローチと、並列スレッドの制御、メモリダンプの方法を解説してきました。

1. コンパイル時に `-g -G -O0` を忘れないこと。
2. カーネルにブレークポイントを張ったら、自動（あるいは手動）で `cuda thread (block, thread)` を切り替えて特定の並列空間をフォーカスすること。
3. デバイスメモリは `x/nf` コマンドでダイレクトに覗き見ること。

この3ステップを武器にするだけで、これまでブラックボックスだったGPU内部の挙動が手に取るようにわかるようになります。「動かないからとりあえず再起動して勘でコードを直す」という泥臭い開発スタイルとは、今日でお別れです。

次回の開発から、ぜひこのプロフェッショナルなデバッグ手法を取り入れてみてください。あなたのコーディングライフが劇的にスムーズになることを、アーキテクトとして確信しています！