DGX SparkでCUDAの範囲外アクセスを特定する

CUDAカーネルは起動しているのに、結果をホストへ戻す段階で初めてエラーが出る。同じプログラムでも、入力を小さくすると動作する。このような場合、まず確認すべきなのはデバイスのメモリアクセスであり、いきなりドライバーを変更することではありません。DGX Sparkでは、小さな独立したプログラムを使い、意図的に範囲外アクセスを発生させ、該当するソース行を特定し、インデックスを修正したうえで、エラー検出が無効になっていないことまで確認できます。モデルの重みは不要で、本番サービスにも触れずに実施できます。

1. 環境を固定し、テストの実施範囲を明確にする

TcpNodeのDGX Sparkは、NVIDIA GB10 Grace Blackwell、ARM64アーキテクチャ、NVIDIA DGX OSを採用し、128GB LPDDR5x統一メモリを搭載しています。統一メモリはCPUとGPUで共有されますが、どのポインターでも割り当て範囲を越えてアクセスできるわけではありません。ここではcudaMallocでデバイスメモリを確保し、有効なアドレス範囲を明確にします。

支払い後、コンソールで接続情報を確認し、SSHでログインします。テストは専用の作業ディレクトリで一般ユーザーとして実行し、リクエストを処理中の推論プロセスにエラーを注入しないでください。

mkdir -p "$HOME/cuda-access-check"
cd "$HOME/cuda-access-check"
uname -m
nvidia-smi
command -v nvcc
command -v compute-sanitizer
nvcc --version
compute-sanitizer --version

アーキテクチャはaarch64と表示されるはずです。ツールのバージョンとドライバー情報を記録し、再検証にも同じ環境を使います。CUDAスタックがプリインストールされていても、すべてのツールが現在のPATHに登録されているとは限りません。コマンドが見つからない場合は、まずインストール済みのツールチェーンとインストール先を確認し、ドライバー一式をいきなり入れ替えないでください。

範囲外アクセスのサンプルは、破棄して構わない独立したプロセスでのみ実行してください。不正なメモリアクセスが発生したら、そのプロセスを終了して実行し直します。同じCUDAコンテキストで得られた後続の結果は、有効な検証材料として扱えません。

2. 変更箇所を1つに絞った範囲外アクセスの再現プログラムを作る

以下をrepro.cuとして保存します。確保する領域は整数1個分だけなので、有効なインデックスはゼロです。コンパイル時にTEST_INDEXで書き込み位置を切り替えます。不具合版と修正版で同じソースを使うことで、入力、割り当てサイズ、実行経路を同時に変更せずに済みます。

#include <cuda_runtime.h>
#include <cstdio>
#include <cstdlib>

static void checked(cudaError_t status) {
    if (status != cudaSuccess) {
        std::fprintf(stderr, "%s\n", cudaGetErrorString(status));
        std::exit(EXIT_FAILURE);
    }
}

__global__ void write_value(int* data) {
    data[TEST_INDEX] = 7;
}

int main() {
    int* data = nullptr;
    int result = 0;
    checked(cudaMalloc(reinterpret_cast<void**>(&data), sizeof(int)));
    checked(cudaMemset(data, 0, sizeof(int)));
    write_value<<<1, 1>>>(data);
    checked(cudaGetLastError());
    checked(cudaDeviceSynchronize());
    checked(cudaMemcpy(&result, data, sizeof(int), cudaMemcpyDeviceToHost));
    checked(cudaFree(data));
    std::printf("result=%d\n", result);
    return result == 7 ? EXIT_SUCCESS : EXIT_FAILURE;
}

2つのプログラムをそれぞれコンパイルします。

nvcc -arch=native -lineinfo -DTEST_INDEX=1 repro.cu -o repro-bad
nvcc -arch=native -lineinfo -DTEST_INDEX=0 repro.cu -o repro-good
sha256sum repro.cu repro-bad repro-good > build.sha256

-arch=nativeは、現在認識されているGPU向けにコンパイルするオプションです。このマシン上での検証には適していますが、異なるデバイスで共通利用するバイナリの生成には使いません。コンパイラーが現在のデバイスに対応していない場合は、まずツールチェーンの互換性を確認し、別のアーキテクチャ向けのオプションを安易に指定しないでください。-lineinfoはデバイス側ソースの行情報を保持します。最初の原因特定では、実行特性を大きく変える完全なデバイスデバッグモードを有効にする必要はありません。

起動時と同期時の両方を確認する理由

cudaGetLastErrorはカーネル起動に関するエラーを確認しますが、カーネルは通常、非同期に実行されます。cudaDeviceSynchronizeは、それ以前に投入されたデバイス処理の完了を待つため、実行中のエラーがここで返されることがあります。両方を個別に確認する必要があり、プログラムが結果を表示したかどうかだけでは判断できません。

3. 検出結果を判定可能な終了コードに反映する

まず不具合版を実行し、標準出力と検出ログを別々に保存します。

if compute-sanitizer --tool memcheck --error-exitcode 86 \
    ./repro-bad > bad.stdout 2> bad.stderr; then
    status=0
else
    status=$?
fi
printf 'exit=%s\n' "$status"
cat bad.stdout bad.stderr

不正なグローバルメモリ書き込みが報告され、write_value内の書き込み処理を特定できるはずです。通常、レポートにはアクセスサイズ、スレッドの位置、割り当て領域とアクセス先アドレスの関係も含まれます。表示形式はツールのバージョンによって異なるため、特定の文言だけを判定基準にしないでください。

--error-exitcode 86は、検出したエラーを自動処理の成否に反映させるためのオプションです。ただし、テスト対象のプログラム自体がすでに失敗している場合は、そのプログラムが返す別の非ゼロ値になることもあります。したがって合格条件は、「終了コードが非ゼロで、レポートが意図した範囲外アクセスを指していること」であり、すべての失敗で86が返ることではありません。

コマンドを通常のログ用パイプにつなぎ、最後のログ処理コマンドの終了ステータスだけを確認する方法は避けてください。上記のリダイレクトと条件分岐なら、実際の実行結果を保持できます。失敗時に即終了する設定のスクリプトでも利用できます。

4. 症状を隠さず、アドレス範囲に基づいて修正する

このサンプルでは整数1個分しか確保していないのに、インデックス1に書き込んでいるため、有効範囲を越えています。実際のプログラムでは、レポートに示されたアドレスを対応する割り当て領域と照合し、インデックス計算、要素の型、ストライド、バッチのオフセットを確認します。

よくある原因には、割り当て時にはバイト数を指定しているのにインデックス側では要素数として扱っていること、最後のスレッドブロックが端数になること、2次元配列の行ストライドと論理的な幅が異なること、解放済みのポインターを別のタスクが再び使っていることなどがあります。割り当てサイズを増やすと一時的にエラーが見えなくなる場合がありますが、それでアルゴリズムの正しさを証明できるわけではありません。

1次元の処理では、通常、アクセス前にindex < countを確認するとともに、インデックス計算に使う整数型で積を表現できることも確認します。共有メモリを使ってスレッドが協調するカーネルでは、一部のスレッドだけを安易に早期リターンさせ、その後で残りのスレッドをブロック単位のバリアに進めてはいけません。新たな同期の問題を引き起こすおそれがあります。

大規模なワークフローでしかエラーが出ない場合は、疑わしい処理段階の直後に一時的な同期チェックを追加し、区間ごとに原因範囲を絞り込めます。チェックポイントは一度に1つだけ追加し、最初のエラーを記録してください。エラー発生後も処理を投入し続け、大量の二次的なエラーを発生させないようにします。

5. 正常系と異常系のテストで修正を検証する

修正版を実行します。

if compute-sanitizer --tool memcheck --error-exitcode 86 \
    ./repro-good > good.stdout 2> good.stderr; then
    status=0
else
    status=$?
fi
printf 'exit=%s\n' "$status"
cat good.stdout good.stderr

この例では、終了コードがゼロであること、検出レポートにメモリアクセスエラーがないこと、出力がresult=7であることの3条件をすべて満たす必要があります。「エラーが出ていない」だけでは不十分です。カーネルが実行されていなかったり、結果が実際には書き込まれていなかったりしても、見かけ上はテストに通ることがあります。

続いて不具合版を再実行し、引き続き失敗することを確認します。この異常系テストにより、検出コマンドが削除されていた、誤って古いバイナリを使っていた、パイプラインが失敗ステータスを握りつぶしていた、といった問題を発見できます。実際のカーネルでは、最小の有効入力、ブロック境界に一致する入力、末尾がブロックに満たない入力もテスト対象に含めます。インターフェースが空入力を許可する場合は、その処理方法も明確にしておきます。

最後に、ソース、コンパイルコマンド、バージョン情報、チェックサム、両方の実行ログを保存します。検出ツールは実行オーバーヘッドを変えるため、検出中の所要時間をそのまま性能評価に使うことはできません。また、memcheckを通過しても、データ競合や同期の問題まで解消されたとは限りません。検証の締めくくりで重要なのは、成功画面を1枚残すことではなく、アクセス範囲の妥当性、計算結果、失敗の伝播について、それぞれ後から確認できる証拠をそろえることです。

よくある質問