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 Single Work-Itemの実行 非線形実行 アウトオブオーダーのループ反復 シリアルリージョン 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.4. Nested Loops

Intel® FPGA SDK for OpenCL™オフライン・コンパイラーループ反復の順序付けのためにパイプライン実行を推測しません。その結果、内部ループの反復回数は、異なるループ反復で異なる可能性があるため、外部ループ反復は、その後の内部ループに対して順序が乱れる可能性があります。

アウトオブオーダーのアウターループ反復の問題を解決するには、アウターループ反復間で変化しない上限と下限を持つ内側ループをデザインします。

Single Work-Itemの実行

FPGA上でのハイスループットのWork-Itemベースのカーネル実行を確実にするために、 Intel® FPGA SDK for OpenCL™オフライン・コンパイラーはある時点で複数のパイプライン・ステージを並列に処理する必要があります。この並列性は、ループの反復をパイプライン化することによって実現されます。

Single Work-Itemでの累積を示す次の簡単なコード例を検討してください。

1 kernel void accum_swg (global int* a, 
                         global int* c, 
                         int size, 
                         int k_size) {
2     int sum[1024];
3     for (int k = 0; k < k_size; ++k) {
4        for (int i = 0; i < size; ++i) {
5            int j = k * size + i;
6            sum[k] += a[j];
7        }
8     }
9     for (int k = 0; k < k_size; ++k) {
10       c[k] = sum[k];
11    }
12 }
各ループ反復の間に、グローバルメモリーaからのデータ値は、 合計[k]に蓄積されます。この例では、4行目の内側ループの開始インターバル値は1で、レイテンシーは11です。外側ループも1以上の開始インターバル値を持ち、レイテンシーは8です。
注: 新しいループの繰り返しは開始間隔（II）と呼ばれます。 IIは、パイプラインが次のループ反復を処理する前に待機しなければならないハードウェア・クロック・サイクルの数を指す。最適にアンロールされたループは、1つのループ反復がクロックサイクルごとに処理されるため、IIの値が1です。
図 43. ループ解析レポート
図 44.  Single Work-Itemカーネルのシステムビュー

次の図は、iの各反復がどのようにブロックに入るかを示しています。

図 45. Inner Loop accum_swg.B2 Execution

外部ループを観測すると、IIの値が1であるということは、スレッドの各反復がすべてのクロックサイクルに入ることもできることを意味します。この例では、 k_sizeが20でsizeが4であると見なされます。これは、最初の8クロック・サイクルでは真であり、外側のループ反復0〜7は、それをストールさせることなく下流に入ることができるからです。スレッド0が内部ループに入ると、それは4回の反復で終了します。スレッド1〜8は内部ループに入ることができず、スレッド0によって4サイクル停止します。スレッド0の反復が完了すると、スレッド1は内部ループに入ります。その結果、スレッド9は、クロックサイクル13で外側ループに入る。スレッド9から20は、 sizeの値である4クロックサイクルごとにループに入る。この例では、外部ループの動的開始間隔が静的に予測される開始間隔1よりも大きく、内部ループのトリップ数の関数であることがわかります。

図 46. Single Work-Itemの実行

非線形実行

ループ構造は線形実行をサポートしていません。次のコード例は、外部ループiに2つの分岐する内部ループが含まれていることを示しています。外側ループの各反復は、1つの内側ループまたは非直線的実行パターンである他方を実行することができます。 

__kernel void structure (__global unsigned* restrict output1,
                         __global unsigned* restrict output2,
                         int N) {
  for (unsigned i = 0; i < N; i++) {
    if ((i & 3) == 0) {
      for (unsigned j = 0; j < N; j++) {
        output1[i+j] = i * j;
      }
    }
    else
    {
      for (unsigned j = 0; j < N; j++) {
        output2[i+j] = i * j;
      }
    }
  }
}

アウトオブオーダーのループ反復

内側ループの反復回数は、外側ループの反復ごとに異なります。次のコード例を検討してください。

__kernel void order( __global unsigned* restrict input,
                     __global unsigned* restrict output
                     int N ) {
  unsigned sum = 0;
  for (unsigned i = 0; i < N; i++) {
    for (unsigned j = 0; j < i; j++) {
      sum += input[i+j];
    }
  }
  output[0] = sum;
}

この例は、 i = 0の場合、内側のループjがゼロ回反復することを示しています。 i = 1の場合、 jは 1回反復します。内部ループの反復回数が変わるため、オフライン・コンパイラーはパイプライン処理を推測できません。

シリアルリージョン

内部ループアクセスが外部ループ依存関係を引き起こすと、ネストされたループ内で直列エリアが発生することがあります。内部ループは、データまたはメモリーの依存性のために、外部ループの反復で直列エリアになります。

定常状態では、外側ループのII =内側ループのII *内側ループのトリップ数となります。 IIが1より大きい内部ループと、直列実行エリアがない外部ループの場合、スレッドを外部ループからインターリーブすることは可能です。

次の式を検討してみましょう。

kernel void serially_execute (global int * restrict A, 
                              global int * restrict B, 
                              global int * restrict result, 
                              unsigned N) {
  int sum = 0;
  for (unsigned i = 0; i < N; i++) {
    int res;
    for (int j = 0; j < N; j++) {
      sum += A[i*N+j];
    }
    sum += B[i];
  }
  *result = sum;
}
この例では、外側ループの依存関係は、内側ループの逐次実行をもたらしました。パフォーマンスの主な違いは、内側ループの定常状態II =内側ループ*（内側ループのトリップ数-1）+レイテンシーです。この例では、内部ループのIIは4のレイテンシーを持ち、IIは外部ループのIIであり、レイテンシーは7です。Nが大きい場合（レイテンシーと比較して400など）、外部ループIIからの影響はほとんどありません。
図 47. カーネルのシステムビュー
図 48. シリアル実行
レベル 2 の表題