Intel® FPGA SDK for OpenCL™: ベスト・プラクティス・ガイド

ダウンロード PDF
ID 683521
日付 12/08/2017
Public
ドキュメント目次
ドキュメント目次 x
1. はじめに 2. カーネルのreport.htmlファイルのレビュー 3. OpenCLカーネルデザインのベスト・プラクティス 4. パフォーマンスのボトルネックを特定するためのカーネルのプロファイリング 5. Single Work-Itemカーネル・パフォーマンスを向上させるための戦略 6. NDRangeカーネルのデータ処理効率を向上させるための戦略 7. メモリーアクセス効率向上のための戦略 8. FPGAエリアの最適化のための戦略 9. 追加情報
1. はじめに x
1.1. FPGA概要 1.2. パイプライン 1.3. シングル・ワーク・アイテム・カーネル対NDRangeカーネル 1.4. マルチ・スレッド・ホスト・アプリケーション
2. カーネルのreport.htmlファイルのレビュー x
2.1. High Level Design レポートの概要 2.2. Report Summary のレビュー 、 2.3. ループ情報のレビュー 2.4. エリア情報の確認 2.5. メモリー・レプリケーションとストールに関する情報の確認 2.6. HTMLレポートの情報に基づいたOpenCLのデザイン例の最適化 2.7. HTMLレポート：エリア・レポート・メッセージ 2.8. HTMLレポート：カーネルデザインの概念
2.3. ループ情報のレビュー x
2.3.1. OpenCLのデザイン例のループ解析レポート 2.3.2. メモリー・アクセス・パターンの例の変更 2.3.3. loop_coalesceを使用してネストループによって消費されるエリアの削減
2.4. エリア情報の確認 x
2.4.1. ソース別エリア分析 2.4.2. システムのエリア分析
2.5. メモリー・レプリケーションとストールに関する情報の確認 x
2.5.1. System Viewerの機能 2.5.2. Kernel Memory Viewer の特長
2.7. HTMLレポート：エリア・レポート・メッセージ x
2.7.1. ボード・インターフェイスのエリア・レポート・メッセージ 2.7.2. 機能オーバヘッドのエリア・レポート・メッセージ 2.7.3. 州のエリア・レポートメッセージ 2.7.4. フィードバックのためのエリア・レポート・メッセージ 2.7.5. 定数メモリーのエリア・レポート・メッセージ 2.7.6. プライベート変数ストレージのエリア・レポート・メッセージ
2.8. HTMLレポート：カーネルデザインの概念 x
2.8.1. Kernels 2.8.2. Global Memory Interconnect 2.8.3. ローカルメモリー 2.8.4. Nested Loops 2.8.5. Single Work-Itemカーネルのループ 2.8.6. チャネル 2.8.7. ロード・ストア・ユニット
3. OpenCLカーネルデザインのベスト・プラクティス x
3.1. データを経由して転送する インテル® FPGA SDK for OpenCL™ チャネルまたはOpenCLパイプ 3.2. ループのアンロール 3.3. 浮動小数点演算の最適化 3.4. アラインメントされたメモリーの割り当て 3.5. 構造体をパディング付きまたはパディングなしで整列する 3.6. ベクトル型要素の類似構造の維持 3.7. ポインター・エイリアシングの回避 3.8. 高価な機能の回避 3.9. Work-ItemID依存の後方分岐の回避
3.1. データを経由して転送する インテル® FPGA SDK for OpenCL™ チャネルまたはOpenCLパイプ x
3.1.1. チャネルとパイプの特性 基本動作 複数のOpenCLカーネルの同時実行 暗黙のカーネル同期 呼び出し間のデータの永続性 課された作業アイテムの注文 3.1.2. チャネルおよびパイプの実行順序 3.1.3. チャネルまたはパイプのバッファー推論の最適化 3.1.4. チャネルとパイプのベスト・プラクティス
3.3. 浮動小数点演算の最適化 x
3.3.1. 浮動小数点表現と固定小数点表現
4. パフォーマンスのボトルネックを特定するためのカーネルのプロファイリング x
4.1. Intel FPGA Dynamic Profiler for OpenCLのベスト・プラクティス 4.2. Intel FPGA Dynamic Profiler for OpenCL GUI 4.3. プロファイリング情報の解釈 4.4. Intel FPGA Dynamic Profiler for OpenCLの制限
4.2. Intel FPGA Dynamic Profiler for OpenCL GUI x
4.2.1. ソース・コード・タブ 4.2.2. カーネル実行タブ 4.2.3. Autorun Capturesタブ
4.2.1. ソース・コード・タブ x
4.2.1.1. ツール・オプション
4.3. プロファイリング情報の解釈 x
4.3.1. ストール、占有、帯域幅 4.3.2. アクティビティー 4.3.3. キャッシュヒット 4.3.4. OpenCLのデザインシナリオ例のプロファイラ解析 4.3.5. 自動プロファイラーデータ
4.3.1. ストール、占有、帯域幅 x
4.3.1.1. チャンネルの停止
4.3.4. OpenCLのデザインシナリオ例のプロファイラ解析 x
4.3.4.1. ハイストール率 4.3.4.2. 低い占有率 4.3.4.3. 低帯域幅効率 4.3.4.4. 高い失業率と高い占有率 4.3.4.5. ノーストール、低占有率、低帯域幅効率 4.3.4.6. ノー・ストール、高占有率、低帯域幅効率 4.3.4.7. 高失速率と低占有率
5. Single Work-Itemカーネル・パフォーマンスを向上させるための戦略 x
5.1. 最適化レポートのフィードバックに基づくてSingle Work-Itemカーネル依存関係のアドレッシング 5.2. メモリー配列へのアクセスによるループキャリー依存関係の削除 5.3. Single Work-Itemカーネルの良いデザイン方法
5.1. 最適化レポートのフィードバックに基づくてSingle Work-Itemカーネル依存関係のアドレッシング x
5.1.1. ループ実行依存関係の削除 5.1.2. Relaxing Loop Carriedの依存関係 5.1.3. Loop Carried依存関係の簡素化 5.1.4. ループで運ばれた依存関係のローカルメモリーへの転送 5.1.5. シフトレジスターの推測によるループキャリー依存関係の削除
6. NDRangeカーネルのデータ処理効率を向上させるための戦略 x
6.1. 最大ワーク・グループ・サイズまたは必要なワーク・グループ・サイズの指定 6.2. カーネルのベクトル化 6.3. 複数のコンピューティング・ユニット 6.4. コンピューティング・ユニット複製とカーネルSIMDベクトル化の組み合わせ 6.5. リソース駆動型最適化 6.6. HTMLレポートのカーネルプロパティとループアンロールステータスの確認
6.2. カーネルのベクトル化 x
6.2.1. スタティック・メモリー統合
6.3. 複数のコンピューティング・ユニット x
6.3.1. コンピューティング・ユニット複製対カーネルSIMDベクトル化
7. メモリーアクセス効率向上のための戦略 x
7.1. メモリーアクセスの最適化に関する一般的なガイドライン 7.2. グローバル・メモリー・アクセスの最適化 7.3. 定数、ローカルまたはプライベート・メモリーを使用したカーネル計算の実行 7.4. ローカルメモリーのバンキングによるカーネル・パフォーマンスの向上 7.5. メモリー・レプリケーションファクタの制御によるローカルメモリーへのアクセスの最適化 7.6. ループパイプラインのメモリー依存性の最小化
7.2. グローバル・メモリー・アクセスの最適化 x
7.2.1. 連続メモリー・アクセス 7.2.2. グローバルメモリーの手動分割
7.2.2. グローバルメモリーの手動分割 x
7.2.2.1. 異種メモリーバッファー
7.3. 定数、ローカルまたはプライベート・メモリーを使用したカーネル計算の実行 x
7.3.1. キャッシュ・メモリー 7.3.2. ローカルメモリーへのデータの事前ロード 7.3.3. プライベート・メモリーに変数と配列の格納
7.4. ローカルメモリーのバンキングによるカーネル・パフォーマンスの向上 x
7.4.1. アレイインデックスに基づくローカル・メモリー・バンクの幾何学的構成の最適化
8. FPGAエリアの最適化のための戦略 x
8.1. コンパイルに関する考慮事項 8.2. ボードバリアントの選択に関する考慮事項 8.3. メモリーアクセスに関する考慮事項 8.4. 算術演算の考慮事項 8.5. データ型の選択に関する考慮事項
9. 追加情報 x
A.1. 改訂履歴
1. はじめに
2. カーネルのreport.htmlファイルのレビュー
3. OpenCLカーネルデザインのベスト・プラクティス
4. パフォーマンスのボトルネックを特定するためのカーネルのプロファイリング
5. Single Work-Itemカーネル・パフォーマンスを向上させるための戦略
6. NDRangeカーネルのデータ処理効率を向上させるための戦略
7. メモリーアクセス効率向上のための戦略
8. FPGAエリアの最適化のための戦略
9. 追加情報

3.1.1. チャネルとパイプの特性

OpenCL™カーネルプログラムにチャネルまたはパイプを実装するには、それぞれの インテル® FPGA SDK for OpenCL™ 特性を考慮してください。

基本動作

チャネルのデフォルト動作はブロックしています。パイプのデフォルト動作はノンブロッキングです。

複数のOpenCLカーネルの同時実行

複数のOpenCLカーネルを同時に実行することができます。同時実行をイネーブルするには、複数のコマンドキューをインスタンス化するようにホストコードを変更します。同時に実行される各カーネルは、別々のコマンド・キューに関連付けられます。

重要:

パイプ固有の考慮事項：

インテル® FPGA SDK for OpenCL™ プログラミング・ガイドの他のOpenCL SDKとの互換性の確保に記載されているOpenCLパイプの変更により、SDKでカーネルを実行することができます。ただし、カーネルのスループットを最大化するわけではありません。 OpenCL仕様バージョン2.0では、カーネルが空のパイプからの読み出しを行わないように、パイプ読み出しの前にパイプ書き込みを行う必要があります。その結果、カーネルは同時に実行できません。 インテル® FPGA SDK for OpenCL™ 同時実行をサポートするため、ホスト・アプリケーションとカーネルプログラムを変更してこの機能を使用することができます。この変更により、アプリケーションのスループットが向上します。ただし、カーネルを別のSDKに移植することはできません。この制限にもかかわらず、変更は最小限であり、両方のタイプのコードを維持するために多大な努力を必要としません。

パイプを含むカーネルの同時実行をイネーブルするには、カーネルコードのdepthの属性をblocking属性（つまり、 __attribute __（（blocking））に置き換えます。 blockingの属性はread_pipeとwrite_pipe関数呼び出しにブロッキング動作を紹介します。コールサイトは、パイプの他端が準備完了になるまで、カーネルの実行をブロックします。

blockingの属性とdepthの属性の両方をカーネルに追加すると、パイプが空のときだけread_pipe呼び出しがブロックされ、パイプがいっぱいになったときにwrite_pipe呼び出しがブロックされます。ブロック動作により、カーネル間の暗黙的な同期が行われ、カーネル同士が互いにロックステップで実行されます。

暗黙のカーネル同期

チャネルをブロックするか、パイプの呼び出しをブロックすることで、カーネルを暗黙的に同期させます。次の例を検討してください。

表 5.  カーネル同期のためのチャネルとパイプコールのブロック
ブロッキング・チャネルコールを持つカーネル ブロッキング・パイプ・コールを持つカーネル 
channel int c0; 

__kernel
void producer (__global int * in_buf) 
{ 
  for (int i = 0; i < 10; i++) 
  {   
    write_channel_intel (c0, in_buf[i]);
  } 
}

__kernel
void consumer (__global int * ret_buf) 
{                 
  for (int i = 0; i < 10; i++) 
  { 
    ret_buf[i] = read_channel_intel(c0);
  }
} 
__kernel
void producer (__global int * in_buf,
  write_only pipe int __attribute__
  ((blocking)) c0) 
{
  for (int i = 0; i < 10; i++) 
{
    write_pipe (c0, &in_buf[i]);
  }
}

__kernel
void consumer (__global int * ret_buf,
  read_only pipe int __attribute__
  ((blocking)) c0) 
{
  for (int i = 0; i < 10; i++) 
  {
    int x;
    read_pipe (c0, &x);
    ret_buf[i] = x;
  }
}

producerカーネルがデータを書き、 consumerカーネルが各ループ反復中にデータを読み込むように、カーネルを同期させることができます。 producerでwrite_channel_intelまたはwrite_pipe呼び出しがread_channel_intelまたはread_pipeコールで任意のデータを書き込まない場合、producerが有効なデータを送信するまで（またはその逆）、consumerはread_channel_intelまたはread_pipeコールをブロックと待機します。

呼び出し間のデータの永続性

write_channel_intelコールがデータをチャネルに書き込んだり、 write_pipe呼び出しがパイプにデータを書き込んだ後も、データはワークグループおよびNDRange呼び出し間で永続的です。Work-Itemがチャネルまたはパイプに書き込むデータは、別のWork-Itemがそこから読み出されるまでそのチャネルまたはパイプに残ります。さらに、チャネルまたはパイプ内のデータの順序は、そのチャネルまたはパイプへの書き込み動作の順序と等価であり、順序は書き込み動作を実行するWork-Itemとは独立しています。

たとえば、複数のWork-Itemがチャネルまたはパイプに同時にアクセスしようとすると、Single Work-ItemだけがそのWork-Itemにアクセスできます。 write_channel_intelコールまたはwrite_pipeコールは、 DATAXという特定のWork-Itemデータをそれぞれチャネルまたはパイプに書き込みます。同様に、チャンネルまたはパイプにアクセスするための最初のWork-Itemは、そこからDATAXを読み出します。読み書き動作のこの順番は、チャネルとパイプをカーネル間でデータを共有するための有効な方法にします。

課された作業アイテムの注文

SDKは、チャネルまたはパイプの読み書き動作の一貫性を維持するためのWork-Itemの順序を強制します。

関連情報
レベル 2 の表題