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

サマリー統計を調査

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

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

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

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

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

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

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

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

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

[Graphics] タブで [Timeline (タイムライン)] ペインに切り替えて、解析の実行中 FPU 使用率を示す [GPU EU Instructions (GPU EU 命令)]メトリックの分散を調査できます。

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

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

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

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

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

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

• 粗粒度バッファー SVM: 共有は、OpenCL* バッファー・メモリー・オブジェクト領域の粒度で行われます。デバイスをまたがるアトミックはサポートされません。

• 細粒度バッファー SVM: 共有は、OpenCL* バッファー・メモリー・オブジェクト領域内のロードとストアの粒度で行われます。デバイスをまたがるアトミックはオプションです。

• 細粒度システム SVM: 共有は、ホストメモリー内で発生するロードとストアの粒度で行われます。デバイスをまたがるアトミックはオプションです。

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

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

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

OpenCL* API (例えば、clWaitForEvents) は、[Thread (スレッド)] エリアにタスクとして示されます。

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

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

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

計算キューを調査

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

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

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

グリッドビューの注目する計算タスクを選択して、ダブルクリックすると [Source]/[Assembly] ウィンドウが開くので、選択したカーネルのコードを解析します。

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

[Graphics] ウィンドウで、計算タスクごとの GPU メトリックデータを調査し、検出した問題に関連する命令を特定するため [Source]/[Assembly] ビューにドリルダウンします。例えば、[Graphics] ウィンドウでサンプラービジーやストールの問題を特定する場合、send 命令を [Assembly] ペインで検索し、これらの命令の頻繁なストールおよびサンプラーの過剰利用を解析します。それぞれの 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]