DGX Spark の CUDA 成果物検証：対象アーキテクチャと PTX

DGX Spark 上で CUDA プロジェクトのビルドは通るのに、カーネル起動時に no kernel image is available for execution on the device が発生する場合、まず確認すべきなのは成果物に含まれる GPU アーキテクチャです。環境一式を再インストールすることではありません。さらに見落としやすいのが、実際には常に PTX の JIT コンパイルで動作しているにもかかわらず、リリース記録ではネイティブのデバイスコードを収録済みとしているケースです。本記事では、独立した小さなプログラムを使い、ターゲットの取得、成果物の検査、二つの実行経路の検証を一つの受け入れ確認手順にまとめます。

1. ホストアーキテクチャと GPU ターゲットを区別する

DGX Spark は NVIDIA GB10 Grace Blackwell を搭載し、ホストアーキテクチャは ARM64 です。OS は Ubuntu 24.04 ベースの NVIDIA DGX OS で、CUDA スタックがプリインストールされています。支払い後、コンソールで接続情報を確認し、SSH で操作できます。

ARM64 は CPU 側の実行ファイルのアーキテクチャを表し、CUDA の Compute Capability（計算能力）は GPU デバイスコードのターゲットを表します。この二つは別々の互換性条件です。file の出力が AArch64 でも、確認できるのはホスト側の形式が正しいことだけであり、使用可能なカーネルが埋め込まれている証拠にはなりません。

SASS はターゲット GPU 向けの機械語コード、PTX はドライバーが JIT コンパイルできる中間表現です。配布パッケージには両方を含められますが、それぞれを個別に検査する必要があります。通常の実行が成功しただけで、両方の経路を検証済みとしてはいけません。

以下の実験では、新しく作成したディレクトリ内にソースコードと実行ファイルを生成するだけで、ドライバー、システムのツールチェーン、業務プログラムは変更しません。ここで得られる結果は互換性の受け入れ確認に用いるものであり、性能測定ではありません。

2. アーキテクチャ番号を推測せず、実機の計算能力を取得する

まず専用の作業ディレクトリを作成し、システムとツールチェーンの情報を保存します。nvidia-smi が表示する CUDA バージョンはドライバーの対応範囲を示すもので、現在使用している nvcc のツールキットバージョンとは限りません。そのため、両方の出力を残します。

mkdir -p "$HOME/cuda-artifact-check"
cd "$HOME/cuda-artifact-check"
uname -m > host-arch.txt
nvidia-smi > driver.txt
nvcc --version > toolkit.txt
nvcc --list-gpu-code > supported-targets.txt

次の内容を probe.cu として保存します。このプログラムは CUDA Runtime を通じてデバイス 0 の計算能力を取得します。失敗時には非ゼロの終了コードを返し、誤って使用されかねないターゲット番号を出力しません。

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

int main() {
    cudaDeviceProp p{};
    cudaError_t e = cudaGetDeviceProperties(&p, 0);
    if (e != cudaSuccess) {
        std::fprintf(stderr, "%s\n", cudaGetErrorString(e));
        return 1;
    }
    std::printf("%d%d\n", p.major, p.minor);
    return 0;
}

nvcc probe.cu -o probe
CC=$(./probe) || exit 1
printf '%s\n' "$CC" > compute-capability.txt
grep -Fx "sm_${CC}" supported-targets.txt || exit 1

ここでは意図的に GB10 のターゲット番号をハードコードしていません。取得した sm_ ターゲットがツールキットの一覧にない場合は、まずツールキットの対応範囲を確認します。適当に古い番号へ置き換えてリリースを進めてはいけません。プローブが失敗した場合は、デバイスへのアクセスやランタイムライブラリの問題を先に解決し、コンパイル段階には進まないでください。

3. SASS と PTX を明示的に生成し、内容を検査する

次の内容を check.cu として保存します。単一のスレッドで整数を書き込み、ホストへコピーして結果を照合するだけのプログラムです。カーネルの起動に失敗しても成功と表示してしまうことがないよう、すべての CUDA 呼び出しでエラーを確認します。

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

#define CUDA_OK(call) do { \
    cudaError_t e = (call); \
    if (e != cudaSuccess) { \
        std::fprintf(stderr, "%s\n", cudaGetErrorString(e)); \
        return 1; \
    } \
} while (0)

__global__ void write_value(int* p) {
    *p = 7;
}

int main() {
    int* device = nullptr;
    int host = 0;
    CUDA_OK(cudaMalloc(&device, sizeof(int)));
    write_value<<<1, 1>>>(device);
    CUDA_OK(cudaGetLastError());
    CUDA_OK(cudaDeviceSynchronize());
    CUDA_OK(cudaMemcpy(&host, device, sizeof(int),
                       cudaMemcpyDeviceToHost));
    CUDA_OK(cudaFree(device));
    if (host != 7) return 2;
    std::puts("PASS");
    return 0;
}

同じ作業ディレクトリで次を実行し、ネイティブコードと PTX の両方を生成するよう明示的に指定します。

CC=$(cat compute-capability.txt)
nvcc -O2 \
  -gencode "arch=compute_${CC},code=[sm_${CC},compute_${CC}]" \
  check.cu -o check

file check > binary-format.txt
cuobjdump --list-elf ./check > elf-list.txt
cuobjdump --dump-sass ./check > sass.txt
cuobjdump --dump-ptx ./check > ptx.txt

ELF 一覧と SASS の出力に該当するターゲットとカーネルが存在するか、PTX の出力に .version、.target、カーネルのエントリーポイントが含まれるかを確認します。ファイルが空、ターゲットが不一致、想定したカーネルがない、といった場合は受け入れ不可とします。コマンドの終了コードを確認するだけでは不十分です。

中間オブジェクトではなく、最終成果物を検査する

実際のプロジェクトには、静的ライブラリのリンク、デバイスリンク、パッケージ化などの工程が入ることがあります。オブジェクトファイルに PTX があることを確認しても、最終的な納品ファイルに残っているとは限りません。最終実行ファイルと実際にロードされる共有ライブラリに対して同じ検査を行い、ファイルのハッシュ値を記録してください。

4. ネイティブ実行と JIT コンパイルを個別に検証する

最初に PTX の JIT コンパイルを禁止し、現在の成果物をネイティブのデバイスコードだけで実行させます。次に PTX 経路を強制し、ディスク上のコンパイルキャッシュも無効にします。どちらも新しいプロセスでテストし、プロセス内ですでにロードされたモジュールが判定に影響しないようにします。

env -u CUDA_FORCE_PTX_JIT \
  CUDA_DISABLE_PTX_JIT=1 \
  ./check > native.log 2>&1
native_rc=$?

env -u CUDA_DISABLE_PTX_JIT \
  CUDA_FORCE_PTX_JIT=1 CUDA_CACHE_DISABLE=1 \
  ./check > ptx-jit.log 2>&1
ptx_rc=$?

printf 'native=%s\nptx_jit=%s\n' \
  "$native_rc" "$ptx_rc" > execution-status.txt

cat execution-status.txt native.log ptx-jit.log
test "$native_rc" -eq 0 && test "$ptx_rc" -eq 0

両方のログに PASS が出力され、終了コードはいずれもゼロである必要があります。テスト用の環境変数は、このコマンドの実行時にだけ適用します。グローバルな環境設定に書き込み、業務プロセスで常時 JIT を強制しないでください。

ネイティブ実行が失敗し、PTX 実行が成功する場合は、まずネイティブターゲットと最終リンク工程を見直します。ネイティブ実行が成功し、PTX 実行が失敗する場合は、PTX が存在するか、ドライバーがそのバージョンに対応しているかを確認します。両方が失敗する場合は、具体的なエラーに基づいてツールチェーン、ランタイムライブラリ、デバイスの状態を調べます。原因の特定をせずにコンパイルオプションを切り替え続けても解決にはなりません。

5. 実験結果をリリース判定の必須条件にする

この小さなプログラムが成功しても、確認できるのは基本的なビルド環境と、このカーネルが現在の環境で動作することだけです。業務アプリケーションでは、拡張ライブラリのロード、入力に応じたカーネルの切り替え、実行時のコード生成などが行われる可能性があります。そのため、実際のエントリーポイントを使ったテストを追加し、主要なデータ型と分岐をカバーする必要があります。

リリースごとに、次の四つの証跡を残すことを推奨します。

• ホストアーキテクチャ、ドライバーの出力、ツールキットのバージョン、実際のコンパイルコマンド。
• GPU の計算能力と、最終ファイルの SASS・PTX 検査結果。
• 二つの実行経路の終了コード、ログ、業務処理の結果を照合した記録。
• ソースコードのバージョン、最終ファイルの SHA-256、動的依存関係の一覧。

sha256sum ./check を使えば、今回の受け入れ確認の対象を特定できます。その後に再リンクした場合は、ファイル名が同じでも再検証が必要です。コンテナーを使用する場合は、実際に配布するイメージ内でツールチェーンと最終ファイルを検査し、実際の実行環境で動作を検証してください。

PTX は互換性を確保するための経路であり、すべてのドライバーでの動作を保証するものではありません。一度の成功から、別のデバイスでも動作すると結論づけることもできません。リリース判定の条件は「現在のドライバー、現在の成果物、検証対象としたエントリーポイントで合格」と明記します。アーキテクチャやツールチェーンが変わった場合は、過去のログを流用せず、証跡を取り直してください。

よくある質問