インテル® VTune™ プロファイラー・ユーザーガイド

インテル® HD グラフィックスとインテル® Iris® グラフィックス上の GPU アプリケーション解析

インテル® VTune™ プロファイラーは、グラフィックス・アプリケーションをプロファイルして、CPU と GPU の両方のアクティビティーを関連付けることができます。

インテル® VTune™ プロファイラーによる GPU 解析は、次の手順に従ってください。

  1. GPU 解析向けにシステムを設定します。

  2. GPU オフロード解析を実行して、アプリケーションが GPU 依存であるか、およびコードが GPU にどれくらい効率良くオフロードされているかを特定します。

  3. DPC++、インテル® メディア SDK と OpenCL* ソフトウェア・テクノロジーを使用する GPU 依存のアプリケーションを詳しく調査するため、GPU 計算/ホットスポット解析を実行します。

カスタム解析を設定して、GPU 利用率データを収集することもできます。これには、解析設定にある [GPU 利用率を解析] オプションを選択します。このオプションは、収集中のオーバーヘッドを軽減します。[プロセッサー・グラフィックス・ハードウェア・イベントを解析] を追加するとオーバーヘッドは中程度、そして [GPU プログラミング API をトレース] オプションを追加すると最大のオーバーヘッドとなります。

GPU 依存のアプリケーションの GPU 利用率を解析

アプリケーションやその処理ステージが GPU 依存であることが判明している場合、[GPU 計算/メディア・ホットスポット] 解析で [特徴付け] オプションを使用して、GPU エンジンが効率良く実行されているか、また改善の余地があるかを調査します。これは、インテル® VTune™ プロファイラーのインテル® グラフィックスのレンダーと GPGPU エンジン向けのハードウェア・メトリックの収集により解析できます。

GPU ハードウェア・メトリックを調査

GPU ハードウェア・メトリックは、GPU アクティビティーを解析して詳細情報を提供し、パフォーマンス改善の可能性を示唆します。インテル® グラフィックス上のレンダーと GPGPU エンジン上で、次のタイプのGPU イベントメトリックを収集するには、GPU 計算/メディア・ホットスポット解析を設定する必要があります: [概要] (デフォルトで収集)、[基本計算 (グローバルとローカル・メモリー・アクセス)]、および [計算拡張] (インテル® Core™ M プロセッサー以降向け)。[解析の設定] で [特徴付け] オプションを選択して必要なグループを指定することで、これらの収集を有効にします。一般的に、イベントグループの [概要] グループで一般的な GPU 実行ユニット、サンプラー、メモリーとキャッシュ・アクセス解析から開始してから、[基本計算 (グローバル/ローカル・メモリー・アクセス)] グループで異なるタイプの GPU メモリーへのアクセス解析を行うことを推奨します。[基本計算] メトリックは、GPU ハードウェア・メトリックと実際の GPU 負荷を関連付けることを可能にする、[GPU 利用率の解析] イベントオプション (デフォルトは GPU 計算/メディア・ホットスポット解析) を有効にして GPU の計算ワークを解析する場合に最も効果的です。また、[完全な計算] イベントグループを選択して、[概要][基本計算] で事前定義されたメトリックを組み合わせて同じビューにそれらを表示することで、GPU 実行ユニットが待機している理由を調査するのに役立ちます。このイベント設定を使用するには、ターゲットの高度な設定で [複数実行を許可] が有効に設定されていることを確認してください。

データが収集されたら、[サマリー] ウィンドウの [EU アレイストール/アイドル] セクションを調査して、実行ユニットが待機している最も典型的な原因を特定します。

設定した事前定義イベントに応じて、インテル® VTune™ プロファイラーは実行ユニットのストール/アイドルに関連するメトリックを解析します。[概要] で事前定義されているデフォルト収集による GPU 計算/メディア・ホットスポット解析には、サンプラービジー、サンプラーがボトルネック、そして GPU L3 帯域幅などの一般的な GPU メモリーアクセスを追跡するメトリックが含まれます。これにより、[EU アレイストール/アイドル] セクションには、サンプラー・ビジー・セクションの GPU 計算タスクのリスト、頻繁にサンプラーをアクセスし GPU L3 帯域幅に依存する最もホットな GPU 計算タスクが表示されます。

サンプラービジー

解析設定で、[基本計算] を選択している場合、インテル® VTune™ プロファイラーは GPU 上の異なるタイプへのデータアクセスを特定するメトリックを解析し、占有率の低い GPU タスクを識別するのに役立つ占有率セクションを表示します。

低い占有率

占有率がアプリケーションの問題であると通知されている場合、タスクが大きすぎるか小さすぎることで、EU アレイアイドルになることが考えられるため、計算タスクのサイズを変更することを検討してください。

[基本計算] 事前定義では、DRAM 帯域幅の解析も有効になります。GPU ワークロードが DRAM 帯域幅依存である場合、対応するメトリック値にフラグが示されます。実行中に DRAM 帯域幅を過度に使用する GPU 計算タスクをテーブルで調査します。

解析設定で、[完全な計算][複数実行を許可] が選択されている場合、インテル® VTune™ プロファイラーは、データ収集に [概要][基本計算] イベントグループの両方を使用し、同じビューに EU アレイストール/アイドルのすべての原因を表示します。

インテル® HD グラフィックスおよびインテル® Iris® グラフィックスのハードウェア・イベントを解析するには、GPU 解析用にシステムがセットアップされていることを確認してください。

時間経過における HW メトリックごとの GPU パフォーマンス・データを解析するには、[グラフィックス] ウィンドウを開き、[タイムライン] ペインに注目します。[グラフィックス] ウィンドウに表示される GPU メトリックのリストは、解析設定で選択されたハードウェア・イベントの事前定義に依存します。

次の例は、GPU 依存のアプリケーションで収集された [概要] メトリックのグループです。

最初に GPU 実行ユニットで次のメトリックを観察します: EU アレイアイドル。アイドルサイクルは無駄なサイクルです。スケジュールされているスレッドがなく、EU の貴重な計算リソースが活用されません。EU アレイアイドルがゼロの場合、GPU は効率良く活用されており、すべての EU にスレッドがスケジュールされています。

ほとんどの場合、最適化では EU アレイストールを最小にし、EU アレイアクティブを最大にします。ただし、メモリー帯域幅に制約されるアルゴリズムは例外です。この場合、(EU アレイアクティブを最大にするのではなく) メモリー帯域幅を特定のプラットフォームのピークに近づけるように最適化します。

メモリーアクセスは、しばしばストールの原因になります。メモリーレイアウトと注意深く設計されたメモリーアクセスの重要性を軽視することはできません。EU アレイストールがゼロでなく GPU L3 ミスに関連しており、アルゴリズムがメモリー帯域幅に制約されていない場合は、メモリーアクセスとレイアウトの最適化を試してみるべきでしょう。

サンプラーアクセスはコストが高く、容易にストールを引き起こします。サンプラーアクセスは、サンプラーがボトルネックサンプラーがビジーのメトリックによって測定できます。

OpenCL* カーネル実行の調査

アプリケーションが、OpenCL ソフトウェア・テクノロジーを使用している場合、[グラフィックス] ウィンドウの [タイムライン] ペインで GPU 計算スレッドのディスパッチ・メトリックを使用して、アプリケーションが GPU 上で多くのワークを計算していることを確認してから、解析を続行しインテル® グラフィックス上で実行されている OpenCL* カーネルの情報を取得します。この解析を行うには、[解析の設定] で [GPU プログラミング API をトレース] オプションを有効にします。GPU 計算/メディア・ホットスポット解析では、デフォルトでこのオプションが有効になっています。

[サマリー] ビューでは、[最もホットな GPU 計算タスク] セクションに GPU 上で実行される OpenCL* カーネルを表示し、パフォーマンス上クリティカルなカーネルを通知します。カーネル名をクリックして、[計算タスク (GPU)/インスタンス] でグループ化された [グラフィックス] ウィンドウを開きます。また、計算タスクでグリッド内のデータをグループ化することもできます。インテル® VTune™ プロファイラーは、次の計算タスクの目的を識別します: 計算 (カーネル)、転送 (OpenCL* ルーチンがホストから GPU へのデータ転送を行います)、および同期 (例えば、clEnqueueBarrierWithWaitList)。

対応するカラムに GPU でのカーネルの実行時間と 1 回の呼び出しの平均時間 (clEnqueueNDRangeKernel の 1 回の呼び出しに対応)、ワークグループのサイズ、カーネルの平均 GPU ハードウェア・メトリックが表示されます。メトリックのカラムヘッダーにカーソルを移動すると、メトリックの説明が表示されます。計算タスクのメトリック値が、インテル® アーキテクチャー向けに設定されているしきい値を超えると、値はピンク色でハイライトされ、パフォーマンス上の問題があることを示します。そのような値にカーソルを移動すると、問題の説明が表示されます。

最初に、実行時間が最も長いホットなカーネルを解析し最適化します。ホットなカーネルには、平均実行時間が長い、あるいは平均実行時間は短くても頻繁に呼び出されるといった特徴があります。どちらの場合も注目すべきです。

OpenCL* カーネルへの転送の詳細を表示してキュー内で費やされた時間を解析するには、[グラフィックス] または [プラットフォーム] ウィンドウの [タイムライン] ペインにある、[計算キュー] データを調査します。

インテル® メディア SDK タスクの実行を調査

[GPU 利用率を解析][GPU プログラミング API をトレース] オプションの両方を有効にして、インテル® メディア SDK プログラム解析を行う場合、[グラフィックス] ウィンドウを使用してインテル® メディア SDK タスクの実行データと GPU ソフトウェア・キューのデータを関連付けます。

コード行ごとの GPU カーネルを解析

コードレベル解析モードで、GPU 計算/メディア・ホットスポット解析を実行して、GPU 解析を GPU オフロード解析で特定されたホットな GPU カーネルに絞り込むことができます。この解析により、パフォーマンス・クリティカルな基本ブロックや、GPU カーネルでのメモリーアクセスに起因する問題を特定して、コード行/アセンブリー命令ごとのパフォーマンス統計が得られます。

関連情報