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

GPU OpenCL* アプリケーション解析

インテル® VTune™ プロファイラーによってアプリケーションが GPU 依存 であり、OpenCL* ソフトウェア・テクノロジーを使用していることが判明している場合、アプリケーションが OpenCL* カーネルを効率良く使用しているか特定するためカスタム解析で [GPU プログラミング API をトレース] 設定オプションを有効にします。デフォルトでは、このオプションは GPU 計算/メディア・ホットスポットと GPU オフロード解析で有効になります。OpenCL* アプリケーションのパフォーマンスを調査するには、GPU 計算/メディア・ホットスポット・ビューポイントを使用します。

次の手順で OpenCL* アプリケーション解析でインテル® VTune™ プロファイラーによって提供されるデータを調査します。

  1. サマリー統計を調査します。

    • GPU 利用率を解析します。

    • 実行ユニット (EU) が、ストールまたはアイドルしている原因を特定します。

    • OpenCL* カーネルによる両方の浮動小数点ユニット (FPU) の過剰使用を特定します。

  2. ホットな GPU OpenCL* カーネルを解析します。

  3. OpenCL* カーネルデータと GPU メトリックを関連付けます。

  4. 計算キューを調査します。

  5. ソースとアセンブリー・コードを解析します。

サマリー統計を調査

アプリケーション・レベルのパフォーマンス特性を提供する [サマリー] ウィンドウからデータ解析を開始します。通常、ベースラインとして、ターゲットの合計実行時間を示す経過時間メトリックに注目します。

このデータをアプリケーションの実行中 GPU エンジンによって使用された GPU 時間と関連付けることができます。

GPU 使用

GPU 時間が経過時間の大部分を占めている場合 (95.6%)、アプリケーションが GPU 依存であることは明白です。ここでは、GPU 時間の 94.4% が OpenCL* カーネルの実行に費やされていることが分かります。

OpenCL* アプリケーションでは、GPU で最も実行時間を費やした OpenCL* カーネルの一覧が示されます。

フラグ付きのカーネルにマウスをホバーすると、実行中に特定されたパフォーマンス上の問題を知ることができます。そのようなカーネル名をリストでクリックすると、計算タスクでグループ化され、合計時間でソートされ、グリッドでそのカーネルが選択された [グラフィックス] ウィンドウが開きます。

解析設定中に使用した事前定義された GPU ハードウェア・イベントによって、インテル® VTune™ プロファイラーは、GPU 実行のストール/アイドルの理由を調査し、サマリーにそれらを表示します。例えば、[基本計算] プリセットでは、GPU L3 帯域幅依存の問題を解析します。

もしくは占有の可能性を示します。

この例では、EU ストールの原因は GPU L3 の高い帯域幅です。リストにある最もホットなカーネルをクリックして、[グラフィックス] ビューに切り替え、キャッシュ再利用の可能性を特定するため、選択したカーネルの [ソース] または [アセンブリー] ビューにドリルダウンします。

アプリケーションの実行で、収集時間の 80% 以上が浮動小数点ユニットの利用に費やされている場合、インテル® VTune™ プロファイラーはそのような値を問題として、FPU を過度に利用するカーネルのリストでハイライトします。

[グラフィックス] タブで [タイムライン] ペインに切り替えて、解析実行中の FPU 利用率を示す GPU EU 命令メトリックの分布を調査できます。

ホットな GPU OpenCL* カーネルの解析

GPU で実行されるすべての OpenCL* カーネルの詳細情報を確認するには、[グラフィックス] ウィンドウに切り替えます。[計算タスクの目的] グループを選択している場合、次の計算タスクのカテゴリーでグループ化されたデータを表示できます: 計算 (カーネル)転送 (OpenCL* ルーチンがホストから GPU へのデータ転送を行います)、および 同期 (例えば、clEnqueueBarrierWithWaitList)。デフォルトでは、グリッドデータは計算タスク/インスタンスでグループ化され、計算タスクのみを示します。OpenCL* 計算タスク以外のプログラム単位で収集されたデータは、[タスク外] エントリーに分類されます。

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

次の例では、Accelerator_Intersect カーネルに、ほとんどの実行時間を費やしています (53.398 秒)。このワークロードで収集された GPU メトリックは、カーネル実行時に高い L3 帯域幅利用率がストールに費やされていることを示しています。計算集約型のコードでは、これはキャッシュ利用率によりパフォーマンスが制限されている可能性を示しています。

グラフィックス・ウィンドウ

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

カーネル・インスタンスが、 OpenCL* 2.0 共有仮想メモリー (SVM) で使用されていると、インテル® VTune™ プロファイラーはそれを識別し、ハードウェアに応じて次のように SVM の使用タイプを表示します。

clCreateKernel の結果は計算カテゴリーの行にあります。2 つの clCreateKernel 呼び出しで (同じソースから) 同じ名前の 2 つの異なるカーネルが作成された場合 (そして、それらが 2 つ以上の clEnqueueNDRangeKernel により実行された場合)、テーブルにはそれぞれのカーネルを表す同じカーネル名の行が 2 つ追加されます。それらが、異なるグローバルまたはローカルサイズ、もしくは異なる SVM 引数セットとして 2 度キューに投入されると、グリッドには個別に表示されます。同じカーネルソースごとにデータを合計するには、[計算タスク/ソース計算タスク (GPU)] のグループ化を使用します。

OpenCL* カーネルデータと GPU メトリックを関連付け

時系列で OpenCL* カーネルの実行を解析するには、[グラフィックス] ウィンドウの [タイムライン] ペイン > [プラットフォーム] タブを調査します。

OpenCL* API (例えば、clWaitForEvents) は、スレッド領域にタスクとして表示されます。

[タイムライン] ペイン: [プラットフォーム] タブ

GPU メトリックと OpenCL* カーネルデータを関連付けます。

GPU ハードウェア・メトリックは、インテル® HD グラフィックス、またはインテル® Iris® グラフィックス向けの [プロセッサー・グラフィックス・イベントを解析] オプションが有効である場合にのみ利用できます。これらのメトリックを収集するには、GPU 解析用にシステムがセットアップされていることを確認してください。

GPU アーキテクチャー・ブロックごとの GPU ハードウェア・メトリックを調査することで、解析する OpenCL* アプリケーションを簡単に検出できます。これを行うには、[グラフィックス] ウィンドウで [計算タスク] グループレベルを選択し、注目する OpenCL* カーネルを選択して [タイムライン] ペインで [IP ブロック分布図] タブをクリックします。インテル® VTune™ プロファイラーは、選択されたカーネルが実行された時間範囲の GPU ハードウェア・メトリックごとのパフォーマンス・データで、プラットフォームのアーキテクチャー分布図を更新します。

現在、この機能は、第 4 世代インテル® Core™ プロセッサーとインテル® Core™ M プロセッサー以降で利用できます。新しい世代のプロセッサーでは、より広いスコープのメトリックが提供されます。

[IP ブロック分布図] を右クリックし、[データを表示] からメトリックデータの表示形式を選択できます。

  • 合計サイズ

  • 帯域幅 (デフォルト)

  • 帯域幅最大値のパーセント

計算キューを調査

OpenCL* カーネル・サブミッションの詳細では、サブミッションと実行の順序に注目し、キューで費やされた時間を解析し、ズームインして、[タイムライン] ペインの [計算キュー] データを調べます。カーネルタスクをクリックすると、上位レイヤーに表示された実行に対応するキュー全体がハイライトされます。同じ名前とサイズのカーネルは、同じ色で示されます。

インテル® VTune™ プロファイラーは、同じ名前とサイズのカーネルを同色で表示します。同期タスクは、影付きの垂直線 で表示されます。データ転送 (ホストシステムから GPU へデータを転送する OpenCL* ルーチン) は、交差斜線 でマークされます。

アタッチモードでは、計算キューがすでに作成されているプロセスにアタッチすると、インテル® VTune™ プロファイラーはこのキューの OpenCL* カーネルのデータを表示しません。

ソースとアセンブリー・コードの解析

グリッドビューで注目する計算タスクを選択してダブルクリックすると、[ソース/アセンブリー] ウィンドウが開き、選択したカーネルのコードを解析できます (ソースが利用できれば)。

コンパイラーが生成した OpenCL* カーネルのアセンブリー・コードを解析し、複雑性を推測し、問題を特定し、クリティカルなアセンブリー行を影響のあるソースコードに関連付け、可能であれば最適化します。例えば、いくつかのコード行が多数のアセンブリー命令にコンパイルされる場合、ソースコードを単純化してアセンブリー行を減らし、コードがよりキャッシュ・フレンドリーになるように検討してください。

[グラフィックス] ウィンドウで、計算タスクごとの GPU メトリックデータを調査し、検出した問題に関連する命令を特定するため [ソース/アセンブリー] ビューにドリルダウンします。例えば、[グラフィックス] ウィンドウでサンプラービジー、またはストールの問題を特定した場合、[アセンブリー] ペインで send 命令を検索して使用頻度を確認します。これらの命令は頻繁にストールを引き起こし、サンプラーを過負荷状態にします。send/sends 命令には、データリード/ライト (型なし/型付きのサーフェスのリードなど)、各種アーキテクチャー・ユニット (サンプラービデオモーション推定) へのアクセス、スレッドの終了 (スレッドの生成元) などを示すコメントが [] で囲まれています。例えば、次の sends 命令は、サンプラーユニットをアクセスするために使用されます。

0x408 260 sends (8|M0) r10:d r100 r8 0x82 0x24A7000 [Sampler, msg-length:1, resp-length:4, header:yes, func-control:27000]

  • -gline-tables-only -s <cl_source_file_name> オプションを使用して中間 SPIR*-V バイナリーが生成されている場合、IL (中間言語) で作成されたカーネルと OpenCL* プログラムでソース/アセンブリー解析が利用できます。

  • #line ディレクティブを使用したソースコードでは、ソース/アセンブリー解析はサポートされません。

  • OpenCL* カーネルがインライン関数を使用する場合、[インラインモード] フィルターバー・オプションを有効にして、グリッド内のインライン関数を表示し、[ソース] ビューで解析できます。

関連情報