インテル® HD グラフィックスとインテル® Iris® グラフィックスを使用するアプリケーションの解析

CPU/GPU の並行性を調査

CPU/GPU 並行性解析を実行して、時間軸における GPU 使用率を調査し、アプリケーションや処理フェーズの一部が CPU もしくは GPU 依存であるかを確認します。これは、インテル® VTune™ Amplifier でサポートされるインテル® グラフィックスと他のサードパーティー製の GPU を搭載したプラットフォーム上で実行されるアプリケーションで利用可能な、最も直感的な解析です。

インテル® VTune™ Amplifier はデータを収集し、解析結果を [GPU Hotspots viewpoint (GPU ホットスポット・ビューポイント)] に表示します。デフォルトで開く [Summary (サマリー)] ウィンドウには、アプリケーションが CPU と GPU をどのように使用しているか高度な統計情報が示され、アプリケーションが GPU 依存であるかどうか理解するのに役立ちます。

次の例は、GPU 依存のアプリケーションの解析結果を示しています。[Summary] ウィンドウで、[Elapsed time (経過時間)] の大部分を [GPU Time (GPU 時間)] が占めていることが分かります。

[Platform (プラットフォーム)] ウィンドウに切り替えて、ソフトウェア・キュー上の GPU 使用率を解析する基本 CPU および GPU メトリックを開き、タイムライン上の CPU 使用率のデータと関連付けます。[Platform] ウィンドウで、GPU が頻繁にビジーになり、ビジー状態と次のビジー状態の間のアイドル時間が非常に短く、GPU ソフトウェアのキューがほとんどゼロにならないことを示している場合、理論的にアプリケーションは GPU に依存しています。ビジー状態の間のギャップが長く、この間 CPU がビジーになる場合、アプリケーションは CPU 依存であると言えます。しかし、このような明確な状況はまれであるため、詳細な解析によりすべての依存関係を明らかにする必要があります。例えば、GPU エンジンがシリアル化されている場合 (GPU エンジンがビデオ処理とレンダリング処理を交互に実行するなど)、誤って GPU 依存と見なされる可能性があります。このケースでは、CPU で実行されるアプリケーションによって、GPU 上で非効率なスケジューリングが生じます。

GPU が時間の経過とともにビジーになる場合、[Graphics (グラフィックス)] ウィンドウに切り替えて、スレッドごとに実行されているワーク (レンダリングや計算) の種類を理解するため詳しく調査します。

[Graphics] ウィンドウにある同じ結果の [Timeline (タイムライン)] ペインでは、[GPU Usage (GPU 使用率)] にギャップがありません。

この例は、レンダーと GPGPU エンジン上のアクティビティー (黄色) に加えて、ブリッター (Blitter) エンジンのアクティビティー (ピンク) を示しています。

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

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

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

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

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

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

解析設定で、[Compute Basic] 事前定義を選択すると、インテル® VTune™ Amplifier は GPU 上の異なるタイプのデータアクセスを特定するメトリックを解析し、占有率の低い GPU タスクの識別に役立つ [Occupancy (占有率)] セクションを表示します。

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

[Compute Basic] 事前定義はまた、DRAM 帯域幅使用率の解析を有効にします。GPU ワークロードが DRAM 帯域幅に依存している場合、対応するメトリック値が示されます。実行中に高い GPU 帯域幅を要求する GPU 計算タスクの一覧を調査できます。

上記の例では、ワークロードの 32.9% の実行時間で DRAM 帯域幅を消費しています。これは、このシステムではパフォーマンスの問題であると考えられます。インテル® VTune™ Amplifier は、テーブルに示す計算タスク (ここでは coalescense_10_1) に注目して、次のような手法を適用することで、メモリー間とのキャッシュライン転送を軽減してデータアクセスを改善することを推奨しています: 1) キャッシュラインが排出される前にすべてのバイトを消費する、2) 計算が制約されるループと帯域幅が制約されるループをマージする。

解析設定で、[Full Compute] 事前定義と 複数回実行モードが選択されている場合、インテル® VTune™ Amplifier は、データ収集に [Overview] と [Compute Basic] イベントグループの両方を使用し、同じビューに EU 配列ストール/アイドルのすべての原因を表示します。

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

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

最初に [GPU Execution Unit (GPU 実行ユニット)] の [EU Array Idle] メトリックを見てみましょう。アイドルサイクルは無駄なサイクルです。スケジュールされているスレッドがなく、EU の貴重な計算リソースが活用されません。[EU Array Idle] がゼロの場合、GPU は効率良く活用されており、すべての EU にスレッドがスケジュールされています。

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

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

サンプラーアクセスはコストが高く、容易にストールを引き起こします。サンプラーアクセスは、[Sampler Is Bottleneck (サンプラーがボトルネック)] と [Sampler Busy] メトリックによって測定できます。

[GPU Hotspots viewpoint] で、[Platform] タブに切り替えて、CPU データ、メモリー帯域幅、割り込み (収集されていれば) などがどのように GPU メトリックデータに関連するか解析します。タイムライン上で対象とする領域を選択して右クリックし、コンテキスト・メニューから [Filter In by Selection (選択でフィルターイン)] を選択して、右ペインでコンテキスト固有の GPU メトリックを調査します。

この例では、選択された範囲で使用されている GPU 実行ユニットごとの統計が表示されています。赤く強調された値は、占有率が低いため GPU 時間の大部分がアイドルであることを示しています。この原因は、非効率なワークのスケジューリングであると考えられます。[GPU Usage] セクションは、GPU エンジンが利用されている場合の GPU エンジンごとの [GPU Time] とアプリケーションの経過時間の比率を表示します。

OpenCL* カーネル実行の調査

アプリケーションが、OpenCL* ソフトウェア・テクノロジーを使用している場合、[Graphics] ウィンドウの [Timeline] ペインで [GPU Computing Threads Dispatch (GPU 計算スレッドのディスパッチ)] メトリックを使用して、アプリケーションが GPU 上で多くのワークを計算していることを確認してから、解析を続行しインテル® グラフィックス上で実行されている OpenCL* カーネルの情報を取得します。この解析を行うには、解析の設定で [Trace OpenCL and Intel Media SDK programs] オプションをオンにします。[GPU Hotspots] 解析では、デフォルトでこのオプションがオンになっています。

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

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

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

OpenCL* カーネルへの転送の詳細を表示し、キュー内で費やされた時間を解析するには、[Graphics] または [Platform] の [Timeline] ペインにある、[Computing Queue (計算キュー)] データを調査します。

インテル® Media SDK タスクの実行を調査 (Linux* のみ)

[Analyze GPU usage] と [Trace OpenCL and Intel Media SDK programs] オプションの両方をオンにして、インテル® Media SDK プログラムの解析を行う場合、[Graphics] ウィンドウを使用してインテル® Media SDK タスクの実行データと GPU ソフトウェア・キューのデータを関連付けます。

[Platform] ウィンドウに切り替えて、[GPU Engine (GPU エンジン)] 領域を調査することで、インテル® Media SDK アプリケーション向けの GPU ソフトウェア・キューと GPU パケット転送の詳細を表示できます。

コードラインごとの GPU カーネルの解析

GPU ホットスポット解析で特定されたホットな GPU カーネルをさらに詳しく調査するため、GPU In-kernel プロファイルを実行できます。この解析では、GPU カーネル内のメモリーアクセスによって生じるパフォーマンス・クリティカルな基本ブロックや問題を特定するのを支援し、コード行やアセンブリー行単位のパフォーマンス特性を提供します。