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. チャネルとパイプの特性 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. 追加情報

2.8.3. ローカルメモリー

ローカルメモリーは複雑なシステムです。異なるレベルのキャッシュがある一般的なGPUアーキテクチャとは異なり、FPGAはローカルメモリーをFPGA内部の専用メモリーブロックに実装します。

ローカルメモリー特性

  • ポート - ローカルメモリーの各バンクには、デザインが同時にアクセスできる書き込みポートと読み出しポートがあります。
  • ダブルポンピング - ダブルポンピング機能により、各ローカル・メモリー・バンクは最大3つのリードポートをサポートします。詳細については、 ダブルポンピングのセクションを参照してください。

ローカルメモリーは複雑なシステムです。異なるレベルのキャッシュがある一般的なGPUアーキテクチャとは異なり、FPGAはローカルメモリーをFPGA内部の専用メモリーブロックに実装します。

図 37. 1つまたは複数のM20Kブロックにおけるローカルメモリーの実装

カーネルのコードでは、 local型の変数としてローカルメモリーを宣言します。 

local int lmem[1024];

Intel® FPGA SDK for OpenCL™ Offline Compilerは、幅、深さ、バンク、レプリケーション、相互接続などのローカルメモリープロパティをカスタマイズします。オフライン・コンパイラーは、コードに基づいてアクセスパターンを分析し、アクセス競合を最小限に抑えるためにローカルメモリーを最適化します。

下の図は、サイズ、幅、深さ、バンク、およびレプリケーションの基本的なローカルメモリープロパティを示しています。

図 38. 複製のないローカルメモリーの例と2回複製された2つのバンク

HTMLレポートでは、ローカルメモリーの全体的な状態は最適であると報告されていますが、複製されており、潜在的に非効率的です。

高効率カーネルをデザインするための鍵は、決してストールしないメモリーアクセスを持つことです。この場合、データパス内のすべての同時メモリーアクセスサイトは、競合することなくメモリーにアクセスすることが保証されています。

複雑なカーネルでは、オフライン・コンパイラーは、メモリーアクセスに競合があるかどうかを推測するのに十分な情報がない可能性があります。その結果、オフライン・コンパイラーはローカルメモリーロードストアユニット（LSU）を推論してメモリーアクセスを調停します。しかし、LSUを推論することは非効率を引​​き起こすかもしれない。詳細については、 ローカルメモリーLSUを参照してください。

図 39. 複雑なローカル・メモリー・システム

オフライン・コンパイラーは、指定した正確なサイズのローカルメモリーを実装するとは限りません。 FPGA RAMブロックは特定のディメンションを持つため、オフライン・コンパイラーはサポートされている次のRAMブロック・ディメンションに切り上げるローカルメモリーサイズを実装します。 RAMブロックの詳細については、デバイス固有の情報を参照してください。

ローカル・メモリー・バンク

ローカル・メモリー・バンクは、デフォルトで最小次元でのみ機能します。複数のバンクを有することにより、同時書込みが可能になる。次の図は、次のローカル変数宣言の実装を示しています。

local int lmem[1024][4];
図 40. Implementation of lmem[1024][4]

ループ内の各ローカルメモリーアクセスには、別々のアドレスがあります。次のコード例では、オフライン・コンパイラーは4つの別々のバンクを作成するためにlmemを推論できます。ループはlmem [] []への4つの同時アクセスを可能にし、最適な構成を実現します。 

kernel void bank_arb_consecutive_multidim (global int* restrict in, 
                                           global int* restrict out) {

  local int lmem[1024][BANK_SIZE];

  int gi = get_global_id(0);
  int gs = get_global_size(0);
  int li = get_local_id(0);
  int ls = get_local_size(0);

  int res = in[gi];

  #pragma unroll
  for (int i = 0; i < BANK_SIZE; i++) {
    lmem[((li+i) & 0x7f)][i] = res + i;
    res >> 1;
  }

  int rdata = 0;
  barrier(CLK_GLOBAL_MEM_FENCE);

  #pragma unroll
  for (int i = 0; i < BANK_SIZE; i++) {
    rdata ^= lmem[((li+i) & 0x7f)][i];
  }

  out[gi] = rdata;

  return;
}
最下位次元でバンクしたくない場合、 bank_bits属性を指定して、メモリーアドレスからバンク選択ビットとして使用するビットを指定します。 bank_bits属性を使用すると、メモリーデータを複数のバンクに分割し、どのアドレスビットを使用してバンクを選択するかを指定できます。指定されたbank_bits属性は、どのメモリーバンクにどのデータ要素が含まれているかを意味します。 
local int [4] [128] __attribute __（（bank_bits（8,7）、bankwidth（4）））;
次の例では、最も低い2つの次元の代わりに7番目と8番目のビットでバンキングが行われます。 
#define BANK_SIZE 4
kernel void bank_arb_consecutive_multidim_origin (global int* restrict in, 
                                                  global int* restrict out) {

local int a[BANK_SIZE][128] __attribute__((bank_bits(8,7),bankwidth(4)));

  int gi = get_global_id(0);
  int li = get_local_id(0);

  int res = in[gi];

  #pragma unroll
  for (int i = 0; i < BANK_SIZE; i++) {
    a[i][((li+i) & 0x7f)] = res + i;
    res >> 1;
  }

  int rdata = 0;
  barrier(CLK_GLOBAL_MEM_FENCE);

  #pragma unroll
  for (int i = 0; i < BANK_SIZE; i++) {
    rdata ^= a[i][((li+i) & 0x7f)];
  }

  out[gi] = rdata;

  return;
}

結果のメモリーのビューは、最初の例の初期ビューと同じです。しかし、バンク・オンに間違ったビットを指定すると、メモリーアービトレーション・ロジックが変化します。

メモリーの次のビューは、メモリーを次のように指定した結果です。 
local int a[4][128] __attribute__((bank_bits(4,3),bankwidth(4)));

コンパイラーがローカルメモリーへのアクセスを別々のアドレスに推論できない場合、ローカルメモリーの相互接続を使用してアクセスを調停し、パフォーマンスを低下させます。

ローカルメモリー複製

ローカルメモリーの複製により、同時に読み出し動作が実行されます。オフライン・コンパイラーは、効率的なローカルメモリーアクセスのためにデザインを最適化して、全体的なパフォーマンスを最大化します。メモリー・レプリケーションは、場合によっては非効率なハードウェアにつながりますが、メモリー・レプリケーションは必ずしもRAMの使用を増加させるとは限りません。

オフライン・コンパイラーが3つ以上のワークグループが同時にローカルメモリーから読み出していることを認識すると、ローカルメモリーをレプリケートします。ローカルメモリーの複製がデザインエリアを大幅に増やす場合、カーネル内の障壁の数を減らすか、またはmax_work_group_size値を大きくして複製のファクタを下げることを検討してください。

ダブルポンピング

デフォルトでは、各ローカル・メモリー・バンクには1つの読み出しポートと1つの書き込みポートがあります。ダブルポンピング機能により、各ローカル・メモリー・バンクは最大3つの読み出しポートをサポートすることができます。

図 41. ローカルメモリーにおけるダブルポンピングのハードウェア・アーキテクチャ

ダブルポンピングを可能にする基本的なメカニズムは、M20Kハードウェアにあります。最初のクロックサイクル中、M20Kブロックはダブルクロックになります。次に、第2のクロックサイクルの間、ポートは多重化されて2つの読み出しポートがさらに形成されます。

ダブルポンピング機能をイネーブルすると、オフライン・コンパイラーはエリア対最大周波数を交換します。オフライン・コンパイラーは、ヒューリスティック・アルゴリズムを使用して最適なメモリー構成を決定します。

ダブルポンピングの利点：

  • 1つの読み出しポートから3つの読み出しポートに増加する
  • RAM使用量を節約する

ダブルポンピングの短所：

  • 冗長ロジックを実装する
  • 最大周波数を下げる可能性がある

次のコード例は、8つの読み出しポートと1つの書き込みポートを持つローカルメモリーの実装を示しています。オフライン・コンパイラーは、ダブルポンピングを可能にし、ローカルメモリーを3回複製して、最大9つの読み出しポートをサポートできるメモリー構成を実装します。

#define NUM_WRITES   1
#define NUM_READS    8
#define NUM_BARRIERS 1

  local int lmem[1024];
  int li = get_local_id(0);

  int res = in[gi];
  #pragma unroll 
  for (int i = 0; i < NUM_WRITES; i++) {  
    lmem[li - i] = res;
    res >>= 1;
  }
  
  // successive barriers are not optimized away
  #pragma unroll
  for (int i = 0; i < NUM_BARRIERS; i++) {
    barrier(CLK_GLOBAL_MEM_FENCE);
  }
  
  res = 0;
  #pragma unroll 
  for (int i = 0; i < NUM_READS; i++) {
    res ^= lmem[li - i];
  }
図 42. Intel FPGA SDK for OpenCL Offline Compilerの8個のリードポートと1個のライトポートを持つlmem []の実装
レベル 2 の表題