NVIDIA OptiX レイトレーシングエンジンは、GPU 上で最適なレイトレーシング性能を実現するためのアプリケーションフレームワークです。OptiX を使用するアプリケーションは、無効な API 引数、真っ黒なフレーム、または数千の同時実行スレッドの下に隠れた GPU 側のバグなど、診断が難しい方法で失敗する可能性があります。

NVIDIA OptiX Toolkit (OTK) のデバッグ機能が役立ちます。OTK は、GPU レイトレーシングアプリケーションで一般的なワークフローをサポートするユーティリティのセットを含む GitHub リポジトリです。OTK は BSD 3 条項スタイルのライセンスで提供されているため、コードを自由にコピーおよび変更できます。

この投稿では、OTK の 2 つのデバッグ機能について説明します。OptiX および CUDA API のエラーコードの一貫したチェックと、ターゲットを絞ったデバイス側デバッグプリントです。OTK には、コンテキストで使用されるデバイス側デバッグプリントを示すサンプルプログラム DemandPbrtScene が含まれています。

OptiX のロギングと検証はどのように機能しますか？

OTK が提供する支援に進む前に、OptiX ログに関する背景知識が役立ちます。OptiX デバイスコンテキストを作成するときに、オプション構造体を提供してログコールバックと検証モードを設定できます。検証モードを OPTIX_DEVICE_CONTEXT_VALIDATION_MODE_ALL に設定すると、OptiX は API 関数への入力を検証できます。 

OptiX は検証エラーのヒューマンリーダブルメッセージをログに書き込みます。ログ出力は、OPTIX_ERROR_INVALID_VALUE のような API エラーを探す最初の場所であるべきです。  検証には API に追加のコストがかかります。デバッグおよびテストビルドではすべての検証を有効にし、リリースビルドでは検証を省略することを推奨します。詳細については、OptiX SDK サンプルを参照してください。

エラーのリターンコードを一貫してチェックする方法

後続の障害が元の問題を隠蔽する前に、できるだけ早くエラーを検出することが最善です。このセクションで詳しく説明します。

API メカニズム

OptiX のほとんどの関数は OptixResult エラーコードを返し、ゼロ以外の場合にエラーを示します。CUDA ランタイム API と CUDA ドライバー API も同様のパターンに従います。3 つの API はすべて以下をサポートします。

  • エラーコード用の個別の列挙型。例: OptixResult
  • エラーコードのシンボリック名を文字列として返す関数。例: OPTIX_ERROR_INVALID_VALUE
  • エラーコードのヒューマンリーダブルエラーメッセージを返す関数。例: Invalid value

これらの関数のシグネチャは 3 つの API でわずかに異なりますが、メカニズムは同じです。

エラーチェックポリシー

すべての API 呼び出しサイトでこれらのエラーコードを手動で処理するのは、面倒でエラーが発生しやすいです。エラーを処理するための一貫したポリシーを適用するには、マクロまたは関数呼び出しを使用する方が良いです。OTK は、エラーが検出されたときに 2 つのポリシーを実装するマクロを提供します。

他のポリシーは、提供されたメカニズムの一部を再利用し、適切なマクロを作成することで簡単に実装できます。

マクロメカニズムの最小限の使用

このエラーチェックメカニズムは、マクロを最小限に使用し、実際の作業を行うためにインライン関数に委譲します。インライン関数の定義にブレークポイントを設定して、関数がエラーを検出したときにデバッガーが実行を停止するようにできます。

マクロは、エラーを引き起こしたコードに関する診断情報を提供するために存在します。

  • expr: マクロに提供された引数の文字列形式。これはエラーコードに評価される式です。
  • __FILE__: マクロが呼び出されたソースファイルの名前。
  • __LINE__: マクロが呼び出されたソースファイル内の行番号。

このマクロ呼び出しサイトからの情報は、実際のエラーチェックを行うインライン関数に渡されます。

API 間の統一されたエラーチェック

インラインのテンプレート関数 checkError は、ステータスコードにエラーがないかチェックし、障害を検出した場合に診断メッセージを作成します。これらの 3 つの API の場合、単純にエラーコードを bool にキャストするだけでエラーを示すのに十分です。3 つの API すべてで、ゼロのステータスコードは成功を示し、ゼロ以外は失敗を示します。

エラーメッセージは次の形式でフォーマットされます。

file(line): expr failed with error nnn (name): message

ここで expr は評価された式、nnn はステータスコードを int にキャストした結果、name はステータスコードのシンボリック名、message はヒューマンリーダブルエラーメッセージです。名前またはメッセージのいずれかが空の場合、関数はそれらを省略します。

インラインのテンプレート関数 makeErrorString は、このメッセージの構築を担当します。結合されたメッセージを構築するために、インラインのテンプレート関数 getErrorName と getErrorMessage を呼び出します。

各 API には API ステータスコード用の個別の型があるため、テンプレート関数を特殊化して、拡張エラー情報を取得する適切な API 呼び出しを実行できます。

使用方法

OTK は、説明したテンプレート関数の必要な特殊化を提供する API ごとに 1 つのヘッダーを提供します（表 1）。

表 1.  エラーチェックヘッダー

使用している API のヘッダーを含め、すべての呼び出しサイトで単一のマクロ OTK_ERROR_CHECK を使用します。次の例では、3 つの API すべてを使用しています。

OTK_ERROR_CHECK( cudaSetDevice( m_deviceIndex ) );
OTK_ERROR_CHECK( cuCtxGetCurrent( &m_cudaContext ) );
OTK_ERROR_CHECK( cuStreamCreate( &m_stream, CU_STREAM_DEFAULT ) );
OTK_ERROR_CHECK( optixInit() );

ターゲットを絞ったデバイス側デバッグプリントを実行する方法 

グラフィックスアプリケーションの問題は、黒い画面をコード化する方法が多すぎることです。

OptiX デバイスコードの問題をデバッグするには、いくつかのアプローチを取ることができます。いくつかの明白な選択肢が思い浮かびます。

  • デバイスコードのデバッグビルドで CUDA デバッガーを使用する
  • printf でデバイスコードのリリースビルドから情報を取得する

多くのアプリケーションは、デバッグモードでコンパイルすると実行が遅くなりすぎて、インタラクティブなデバッガーの使用が妨げられます。

printf スタイルのデバッグの主な難点は、GPU 上で同時に実行されるスレッドが非常に多いため、出力の洪水に溺れてしまう可能性があることです。また、問題がアプリケーションとの一定の相互作用の後にのみ発生する可能性もあります。  問題が視覚的に現れる前のデバッグ出力は、目的の情報を見つけるのを妨げるノイズに過ぎません。

DebugLocation

ヘッダー <OptiXToolkit/ShaderUtil/DebugLocation.h> は、デバッグ出力のための再利用可能なメカニズムを提供します。構造体 DebugLocation が動作を制御します。

struct DebugLocation
{
    bool  enabled;
    bool  dumpSuppressed;
    bool  debugIndexSet;
    uint3 debugIndex;
};

enabled メンバーは、メカニズム全体をオンまたはオフにします。dumpSuppressed メンバーは、メカニズムが有効な場合でもデバッグ出力をオフにします。debugIndexSet メンバーは、有効な起動インデックスが debugIndex に格納されていることを示します。

メカニズムは、以下の条件が真の場合にデバッグ情報を出力します。

  • enabled が true
  • dumpSuppressed が false
  • debugIndexSet が true、そして
  • 現在の起動インデックスが debugIndex と一致する

デバッグ出力のインタラクティブな制御のために、OptiX パイプラインの起動パラメータに DebugLocation 構造体のインスタンスを含めます。

debugInfoDump 関数

テンプレート関数 debugInfoDump は、デバッグ情報を出力するためのインターフェースを提供します。

template <typename Callback>
static __forceinline__ __device__
bool debugInfoDump( const DebugLocation& debug,
                    const Callback &callback )

Callback テンプレートパラメータは、以下に一致する struct または class である必要があります。

struct Callback
{
    void setColor( float red, float green, float blue );
    void dump( const uint3& index );
};

setColor メソッドは、デバッグ場所を簡単に識別できるように、視覚的なボックスを描画するために使用されます。典型的な使用法は、現在の起動インデックスに対応する出力ピクセルの色を設定することです。デバッグ場所の視覚的な表示が不要な場合は、メソッドを単に空にすることができます。

dump メソッドは、提供された起動インデックスでアプリケーションが関連すると考える情報を出力するために使用されます。

デバッグ場所の表示

有効な場合、コールバック構造体の setColor メソッドは、デバッグ場所を示すために画面にボックスを描画します。ダンプ出力が抑制されている場合でも表示されます。ボックスは無効になっている場合は非表示になります。

デバッグ場所の赤いピクセルは、1 ピクセル幅の黒いボックスの内側にあり、それ自体が 1 ピクセル幅の白いボックスの内側にあります。これにより、ダンプメッセージが発生している場所の高コントラストインジケーターが提供されます。出力バッファが従来のカラーバッファでない場合は、提供された赤、緑、青の値を視覚化のための特徴的な値にマッピングできます。

ワンショットモード

デバッグ出力に溺れるのを避けるために、ワンショットモードでデバッグ出力を保持することが役立ちます。ここで、出力はユーザー制御に応じて一度ダンプされます。このモードを次のようにシーケンスできます。

  1. DebugLocation メカニズムを有効にする
  2. 通常どおり起動する
  3. ユーザーがインタラクティブにデバッグ場所を選択すると、dumpSuppressed を true に、debugIndexSet を true に、debugIndex を選択した場所に設定する
  4. 後続の起動はデバッグ場所を表示しますが、メカニズムはダンプを提供しません
  5. ユーザーは、アプリケーションと対話して適切な状態に操作し、デバッグ場所を移動する可能性があります
  6. ユーザーが現在の場所のデバッグ情報を要求すると、dumpSuppressed を false に設定する
  7. デバッグ出力を取得するために起動する
  8. 起動後に dumpSuppressed を true に戻す

DemandPbrtScene の例

OTK の DemandPbrtScene の例 は、pbrt version 3 シーン の需要ロードジオメトリを示しています。ワンショット動作とデバッグ出力のインタラクティブな切り替えおよびデバッグ位置のインタラクティブな選択を含む DebugLocation メカニズムを使用しています。UI フレームワークとして ImGui を使用しています。

デバッグ出力のコントロールと画像の右側で強調表示されたデバッグピクセルを示すスクリーンショット。
図 1. 強調表示されたデバッグピクセル付きの DemandPbrtScene デバッグコントロール

OTK からこの例を実行するには、pbrt-v3 のシーンファイル が必要です。

OptiX Toolkit は、一般的な OptiX 開発の問題に対する再利用可能なデバッグおよびテストユーティリティを提供します。一貫した API エラーチェックとターゲットを絞ったデバイス側デバッグ出力です。コードは NVIDIA/optix-toolkit GitHub リポジトリから入手できます。

始める準備はできましたか？GitHub から OptiX Toolkit をダウンロードし、デバッグビルドで OptiX 検証を有効にし、CUDA および OptiX 呼び出しを OTK_ERROR_CHECK でラップし、DebugLocation を使用して開発の早い段階で GPU 側のバグを分離することから始めましょう。OTK は寛容な BSD 3 条項スタイルのライセンスで利用できるため、これらのユーティリティを独自の OptiX アプリケーションに直接コピー、適応、および統合できます。