インテル® SDK for OpenCL* Application
インテル® VTune™ プロファイラー
OpenCL
最適化
解析

インテル® VTune™ Amplifier XE: インテル® HD グラフィック上の OpenCL* パフォーマンスの解析を始めましょう

同カテゴリーの次の記事

インテル® Xeon Phi™ コプロセッサー上でインテル® VTune™ Amplifier XE 2013 を利用して OpenCL* アプリケーションを解析する方法

この記事は、インテル® デベロッパー・ゾーンに掲載されている「Intel® VTune™ Amplifier XE: Getting started with OpenCL* performance analysis on Intel® HD Graphics」 (http://software.intel.com/en-us/articles/intel-vtune-amplifier-xe-getting-started-with-opencl-performance-analysis-on-intel-hd-graphics) の日本語参考訳です。

はじめに

インテル® HD グラフィックス (この記事では GPU とします) を使用してレンダリング、ビデオ処理、計算を行うアプリケーションをプロファイルする場合、インテル® VTune™ Amplifier XE が役立ちます。インテル® VTune™ Amplifier XE は、GPU 全体のアクティビティー (グラフィックス処理、メディア処理、計算) を監視し、インテル® インテグレーテッド・グラフィックスのハードウェア・メトリックを収集して、GPU 上の OpenCL* アクティビティーの詳細を提供します。そして、それらを CPU プロセスとスレッドに関連付けて表示します。図 1 にインテル® VTune™ Amplifier XE で表示される新しい情報を示します。

図 1. メディアと OpenCL* の相互運用性サンプル・アプリケーションを第 3 世代インテル® Core™ プロセッサーで実行した場合の CPU/GPU アクティビティーの関連付けと GPU OpenCL* カーネルの情報を表示。

GPU アクティビティーを追跡することで、次のことが分かります。

  • ワークロードが CPU バインドか GPU バインドか
  • GPU で実行するホットな OpenCL* カーネルがあるかどうか
  • インテル® インテグレーテッド・グラフィックスが最大限に活用されているか。されていない場合、どの程度の余力があるか。

アプリケーションのチューニングは、CPU と GPU のそれぞれの最適化と、両者間の通信の最適化を対象とします。ここでは、GPU 側にのみ注目し、GPU OpenCL* カーネルの解析と最適化について説明します。

この記事は、インテル® VTune™ Amplifier XE とそのワークフロー、主な概念を事前に理解していることを前提にしています。これらについての詳細は、インテル® VTune™ Amplifier XE ヘルプを参照してください。

GPU 解析は、現在も活発に研究が行われている分野です。インテル® VTune™ Amplifier XE チームでは、皆様からのご意見、ご要望をお待ちしております。http://software.intel.com/en-us/forums/topic/380102 (英語) からぜひご意見をお寄せください。

現在、この機能は Windows* でのみサポートされています。GPU の使用状況に関する一般的なデータは、すべての Windows* 7 または Windows* 8 システムで利用できますが、インテル® インテグレーテッド・グラフィックスのハードウェア・メトリックと GPU OpenCL* データは、第 3 世代および第 4 世代インテル® Core™ プロセッサーでのみ利用可能です。

パフォーマンス・データ収集の設定

重要: ハードウェア・アクセラレーションを利用してグラフィックス・アプリケーションを実行する場合と同様に、リモート・デスクトップ接続により GPU データ収集を行うことはできません。GPU データ収集を実行するには、インテル® VTune™ Amplifier XE をターゲット・コンピューターのコンソールから実行するか、VNC を介してターゲット・コンピューターにアクセスする必要があります。

一般的な GPU の使用状況を監視するには、管理者としてインテル® VTune™ Amplifier XE を実行します。

[New Analysis (新しい解析)] ダイアログで、[Advanced Hotpots (詳細 hotpot)] 解析を選択して GPU アクティビティーを監視します。GPU 収集は 2 つのコントロールによって制御されます: [Analyze DirectX pipeline events (DirectX パイプライン・イベントの解析)] と [Analyze Processor Graphics hardware events (プロセッサー・グラフィックス・ハードウェア・イベントの解析)] (図 2)。

[Analyze DirectX pipeline events (DirectX パイプライン・イベントの解析)] をオンにし、[Analyze Processor Graphics hardware events (プロセッサー・グラフィックス・ハードウェア・イベントの解析)] を [Overview (全般)] にすることを推奨します。

図 2. GPU 収集を有効にする設定

GPU で OpenCL* のプロファイルを有効にするには、[Trace OpenCL kernels on Processor Graphics (プロセッサー・グラフィックスで OpenCL カーネルを追跡する)] をオンにします。

収集と後処理が完了したら、図 1 に示す [Graphics (グラフィックス)] タブに切り替えて、GPU アクティビティーの詳細、CPU プロセスおよびスレッドとの関連性を確認します。

パフォーマンス・データの解析

この記事では、GPU OpenCL* アプリケーションの解析に注目します。OpenCL* アプリケーションの場合、GPU パフォーマンスと最適化の可能性を解析するには、OpenCL* の概念と GPU アーキテクチャーへのマップ方法を理解する必要があります。詳細は、『Intel® SDK for OpenCL* Applications – OpenCL* Optimization Guide』(http://software.intel.com/sites/products/documentation/ioclsdk/2013/OG/index.htm) を参照してください。以下のセクションでは、インテル® インテグレーテッド・グラフィックス・アーキテクチャーの基本、および GPU メトリックと OpenCL* プリミティブ間の関係について説明します。

ホットな GPU OpenCL* カーネル

図 1 のように、[Grouping (グループ)] を [Computing Tasks (GPU) (コンピューティング・タスク (GPU))] に切り替えると、GPU で実行しているすべての OpenCL* カーネルの情報が得られます。

対応する列に GPU でのカーネルの実行時間と 1 回の呼び出しの平均時間 (clEnqueueNDRangeKernel の 1 回の呼び出しに対応)、ワークグループのサイズ、カーネルの平均 GPU ハードウェア・メトリックが表示されます。

各 clCreateKernel 呼び出しにより、[Source Computing Task (GPU) (ソース・コンピューティング・タスク (GPU))] テーブルに行が追加されます。2 つの clCreateKernel 呼び出しでは、(同じソースから) 同じ名前の 2 つの異なるカーネルが作成された場合 (そして、それらが 2 つ以上の clEnqueueNDRangeKernel 呼び出しにより実行された場合)、テーブルにはそれぞれのカーネルを表す同じカーネル名の行が 2 つ追加されます。[Computing Task (GPU) (コンピューティング・タスク (GPU))] は、同じカーネルのワークグループのローカル/グローバルサイズが等しいインスタンスをグループ化するので、ワークグループ・サイズが異なる 1 つのカーネルを実行する場合、それぞれのインスタンスは、[Computing Task (GPU) (コンピューティング・タスク (GPU))] グループで個別の行になります。

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

GPU アーキテクチャーについて

GPU とは、小さなコアの配列 (実行ユニット (EU)) で計算処理が行われる高度な並列マシンです。各 EU は、複数の軽量なスレッドを同時に実行します。スレッドの 1 つが実行されると、ほかのスレッドがメモリーなどからのデータを待機するためストールしていても、そのストールを隠蔽することができます。

GPU の性能を最大限に利用するため、並列アプリケーションはできるだけ多くのスレッドをスケジュールして、アイドルサイクルを最小限に抑えようとします。ストールを最小限に抑えることは、GPU アプリケーションでも非常に重要です。

図 3. GPU 機能ユニットに関連したインテル® インテグレーテッド・グラフィックスのハードウェア・メトリックの概略図

インテルの GPU ハードウェア・メトリックは、サンプリング期間の GPU リソースの使用状況に関する情報を提供します。例えば、EU がアイドル状態、ストール状態、アクティブ状態であったサイクルの比率や、メモリーアクセスとほかの機能ユニットに関する統計情報などが分かります。図 3 は、インテル® VTune™ Amplifier XE で表示可能なインテル® GPU のさまざまな部分で収集されるメトリックの概略図です。

各メトリックの意味

GPU ハードウェア・メトリックは、GPU ハードウェア・リソースが効率良く使用されているか、そしてパフォーマンス向上の可能性があるかどうかを示します。

メトリックの多くは、サンプリング期間中のすべてのサイクルに対して GPU 機能ユニットが特定の状態であったサイクルの比率で表されます。以下に、3 つの主要なメトリックを示します。

EU アクティビティーを示す EU Array Active (EU 配列アクティブ) メトリック:

EU ストールを示す EU Array Stalled (EU 配列ストール) メトリック:

EU アイドルを示す EU Array Idle (EU 配列アイドル) メトリック:

この 3 つのメトリックの値は 0 から 1 の範囲内になります。

メトリックを利用した最適化

最初に EU Array Idle (EU 配列アイドル) メトリックを見てみましょう。アイドルサイクルは無駄なサイクルです。スケジュールされているスレッドがなく、EU の貴重な計算リソースが活用されません。EU Array Idle (EU 配列アイドル) がゼロの場合、GPU は効率良く活用されており、すべての EU にスレッドがスケジュールされています。ゼロでない場合、スレッドのスケジュールに問題があるか (例えば、並列に実行できるワークグループの数が足りないなど)、あるいはワークグループ内のスレッド間でインバランスが生じていることが考えられます。インバランスは、ワークグループ内のいくつかのスレッドがすでにワークを完了しているのに、ほかのスレッドがまだ実行している場合に発生します。すべてのスレッドが完了するまで、別のワークグループは開始できません。

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

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

1 つの最適化手法は、共有ローカルメモリー (SLM) を使用することです。グローバルバッファーのコンテンツを SLM にコピーすることで、ワークグループ内のワークで共同で使用できます。SLM を使用する場合は、大きなSLM 領域を要求すると利用可能な L3 領域が減り、並列に実行できるワークグループの数も制限されることに留意してください。

次のように、カーネルソースで SLM の一部を割り当ててみると良いでしょう。

local float temp[BUFF_SIZE];

4*BUFF_SIZE は、使用する SLM のサイズです。これで、カーネルの平均時間とメトリックが大きく変わるかどうかを確認します。変わらない場合、指定した SLM サイズは、カーネルの L3 領域と並列に実行するワークグループの数に大きく影響しないため、SLM ベースの最適化を利用できます。

カーネルの SLM アクセスは、GPU Shared Local Memory Read/Write (GPU 共有ローカルメモリーの読み込み/書き込み) メトリックで評価されます。

現在、インテル® VTune™ Amplifier XE は CPU/GPU メモリー転送に関する情報を提供していません。この機能は、将来のリリースで追加される予定です。メモリー転送と転送時間を把握することは重要です。データのコピーを回避するには、できるだけ CPU 上のアライメントされたバッファーと CL_MEM_USE_HOST_PTR を使用します。

サンプラーアクセスはコストが高く、容易にストールを引き起こします。これは、カーネル内の read_image<> 呼び出しによって引き起こされ、Sampler Is Bottleneck (サンプラー・ボトルネック) および Sampler Busy (サンプラービジー) メトリックに影響します。サンプラーアクセスはできるだけ避けて、read_image<> 呼び出しの代わりに単純なメモリーバッファーを使用すると良いでしょう。ただし、これは常に可能というわけではありません。Sampler Is Bottleneck (サンプラー・ボトルネック) がゼロでない場合、ストールを減らす最良の方法は、隣接するワークにおいて read_image<> 呼び出しで隣接するピクセルを要求することです。隣接するピクセルにより高い局所性が保証され、サンプラーは短時間でデータを返すことができます。

インテル® VTune™ Amplifier XE のメトリックの一覧と詳細は、付録を参照してください。

まとめ

この記事では、GPU OpenCL* カーネルの解析と最適化を行うインテル® VTune™ Amplifier XE 機能について紹介しました。インテル® VTune™ Amplifier XE チームでは、これらの機能について、皆様からのご意見、ご要望をお待ちしております。http://software.intel.com/en-us/forums/topic/380102 (英語) からぜひご意見をお寄せください。

インテル® SDK for OpenCL* Applications 2013 について

インテル® SDK for OpenCL* Applications 製品の最新バージョンであるインテル® SDK for OpenCL* Applications 2013 は、OpenCL* 1.2 対応の第 3 世代および第 4 世代インテル® Core™ プロセッサー・ベースの Windows* 7/Windows* 8 システムで実行する OpenCL* アプリケーション向け統合ソフトウェア開発環境です。OpenCL* アプリケーション開発のビルド、デバッグ、チューニング・プロセスを支援するツールを開発者に提供します。効率良く開発を行えるように、インテル®VTune™Amplifier XE とともに使用することを推奨します。

インテル® SDK for OpenCL* Applications 2013 は、intel.com/software/opencl (英語) から無料で利用できます。

用語集

GPU – グラフィックス・プロセシング・ユニット。

EU – 実行ユニット。GPU 上で計算処理を実行する配列の 1 つのコア。

SLM – 共有ローカルメモリー。1 つの OpenCL* ワークグループのワークで共有される、ソフトウェアにより制御されるメモリー。

LLC – 最終レベルキャッシュ。複数の CPU コアと統合 GPU で共有される CPU アンコアにあるキャッシュ。

付録
インテル® インテグレーテッド・グラフィックス・ハードウェア向けのインテル® VTune™ Amplifier XE メトリック

メトリック
プリセット 1、2		 説明 使用方法

EU Array Active (EU 配列アクティブ)

すべてのコアで命令の実行に費やされたすべてのサイクルの正規化された合計。

計算により制約されるコードの場合、1 に近くなります。

EU Array Stalled (EU 配列ストール)

すべてのコアでストールに費やされたすべてのサイクルの正規化された合計。少なくとも 1 つのスレッドがロードされているが、コアがストールされている状態。

ゼロでない場合は、ストールの原因を調査します (例: メモリーまたはサンプラーアクセス)。

EU Array Idle (EU 配列アイドル)

コアでスレッドがスケジュールされなかったすべてのサイクル (すべてのコア分) の正規化された合計。(1 – (EU Active + EU Stalled)) と等しい。

インバランスやスレッドのスケジュール問題を見つけます。理想的な値はゼロ です。

Compute Shader Threads (計算シェーダースレッド)

計算処理を行うためすべての EU で開始されたスレッドの数。

計算処理とグラフィックス処理を区別するのに役立ちます。純粋なグラフィックス・アクティビティー (計算シェーダーなし) の場合、この値はゼロになります。
メトリック
プリセット 1		 説明 使用方法

GPU Memory Reads/Writes (GPU メモリーの読み込み/書き込み)

GPU とチップのアンコア (LLC) およびメモリー間の読み込み/書き込み。内部 GPU L3 キャッシュミスになり、アンコアまたはメインメモリーから読み込まれたすべてのメモリーアクセス。

アンコア (LLC) またはメインメモリーから読み込む場合、レイテンシーが長いため、よくストールが発生します (EU Stalled != 0 で確認できます)。
カーネルに十分な計算処理がある場合、このレイテンシーは隠蔽されストールは発生しません (EU Stalled ~ 0)。
メモリーアクセスの最適化は、EU Stalled !=0 で L3 ミスが多数発生している場合のみ考慮すべきです。計算とメモリー・レイテンシーにより制約されるコードの場合は、メモリーのレイアウトとアクセスを改善してください (例: SLM の使用やL3 データの再利用)。メモリー帯域幅により制約されるコードの場合は、EU ストールの有無に関係なく、最大帯域幅を達成できるように最適化してください (例: 1 つのワークの作業量を増やしてワークグループの数を減らすなど)。

L3 Cache Misses (L3 キャッシュミス)

GPU L3 キャッシュのすべての読み込み/書き込みミス。

Sampler Busy (サンプラービジー)

すべてのコアでサンプラーが (read_image などにより) ビジー状態であったすべてのサイクルの正規化された合計。

およそ 1 の場合、ストールが原因である可能性があります。サンプラーの使用を減らしてください (例: clCreateBuffer で作成した単純なバッファーを利用します)。

SamplerIsBottleneck (サンプラー・ボトルネック)

入力用の FIFO キューが一杯のため、サンプラーが EU をストールし、出力用の FIFO キューが空になってしまった状態 (EU はサンプラーへの要求の送信を待機しなければならない)。

およそ 0.01 よりも大きい場合、ストールが原因である可能性があります。サンプラーの使用を減らすか、サンプラーへのアクセスの局所性を高めます。

Texture Read (テクスチャーの読み込み)

サンプラーキャッシュのサンプラー・ユニット・ミス。

ストールの原因である可能性があります。サンプラーの使用を減らすか、サンプラーへのアクセスの局所性を高めます。
メトリック
プリセット 2		 説明 使用方法

Untyped Memory Reads/Writes (型指定のないメモリー読み込み/書き込み)

clCreateBuffer で作成されたバッファーへのメモリーアクセス。

すべてのアクセスをカウントします。

Typed Memory Reads/Writes (型指定のあるメモリー読み込み/書き込み)

型指定のあるバッファーへのメモリーアクセス (例: clCreateImage で作成されたバッファーへの書き込み)。ただし、イメージの読み込みは、サンプラーアクセスとテクスチャー読み込みでカウントされる。

すべてのアクセスをカウントします。

SLM Reads/Writes (SLM 読み込み/書き込み)

共有ローカルメモリーへのメモリーアクセス。

SLM へのすべてのアクセスをカウントし、SLM 帯域幅を報告します。SLM が理論的なピーク帯域幅とどの程度近いかを確認できます。

本資料に含まれるソフトウェア･ソース･コードはソフトウェア･ライセンス契約に基づいて提供されるものであり、その使用および複製はライセンス契約で定められた条件下でのみ許可されます。

Intel、インテル、Intel ロゴ、Intel Core、VTune は、アメリカ合衆国およびその他の国における Intel Corporation の商標です。

© 2013 Intel Corporation. 無断での引用、転載を禁じます。

* その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

OpenCL および OpenCL ロゴは、Apple Inc. の商標であり、Khronos の使用許諾を受けて使用しています。

コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

関連記事