OpenCL Specification version 1.0、Khronos Group™によるOpenCL Specification version 1.0で説明されているように、OpenCLの概念とアプリケーション・プログラミング・インターフェイス（API）に精通していることを前提としています。 また、OpenCLアプリケーションの作成経験があることを前提としています。

FPGA用のOpenCLアプリケーションの最高のパフォーマンスを実現するには、基礎となるハードウェアの詳細に慣れてください。 さらに、OpenCLアプリケーションをFPGAに変換してマッピングするコンパイラーの最適化についても理解してください。

OpenCL Specification version 1.0の詳細については、Khronos Group ウェブサイトのOpenCL Reference Pagesを参照してください。 OpenCL APIとプログラミング言語の詳細については、OpenCL Specification version 1.0を参照してください。

FPGAはいくつかの小さなコンピューティング・ユニットで構成されています。次の図に示すように、コンピューティング・ユニットをプログラミングして接続することで、カスタムデータパスをファブリックに直接組み込むことができます。データフローはアーキテクチャに直接プログラミングされています。

図 1. FPGAアーキテクチャ

FPGAでは、ビットマスキング、シフト、加算などの低レベルの動作がすべて設定可能です。また、これらの動作を任意の順序でアセンブルできます。計算パイプラインを実装するために、FPGAはルックアップ・テーブル（LUT）、レジスター、オンチップメモリー​​、および算術ハードウェア（例えば、デジタル信号プロセッサー（DSP）ブロック）の組み合わせを再構成可能な接続のネットワークを介して統合します。その結果、FPGAは高度なプログラマビリティを実現します。 LUTは、さまざまなロジック機能を実装します。例えば、LUTをリプログラミングすることにより、ビット単位のAND論理関数からビット単位のXOR論理関数に演算を変更することができます。

アルゴリズムを高速化するためにFPGAを使用することの主な利点は、幅広く異機種でユニークなパイプライン実装をサポートすることです。この特性は、対称型マルチプロセッサー、DSP、およびグラフィックス処理ユニット（GPU）などの多くの異なるタイプの処理ユニットとは対照的です。これらのタイプのデバイスでは、同じ汎用コンピューティング・ハードウェアを複数回複製することによって並列性が実現されます。しかし、FPGAでは、アルゴリズムが実行するロジックのみを複製することで並列処理を実現できます。

プロセッサーは、各クロックサイクルで実行できる作業量を制限する命令セットを実装します。たとえば、ほとんどのプロセッサーには、次のCコードを実行できる専用命令がありません。

E = (((A + B) ^ C) & D) >> 2;

このCコード例のための専用の命令がないと、CPU、DSP、またはGPUは、動作を実行するために複数の命令を実行する必要があります。対照的に、FPGAは、ソフトウェア・アルゴリズムに必要な命令セットを実装できるハードウェア・プラットフォームと考えることができます。上記のコード例を1クロックサイクルで実行する一連の動作を実行するようにFPGAをコンフィグレーションすることができます。 FPGA実装は、特殊な追加ハードウェアとビット単位のXORおよびAND演算を実行するLUTを接続します。その後、デバイスはプログラマブルな接続を使用して、ハードウェア・リソースを消費することなく2ビット右シフトを実行します。この動作の結果は、複雑なパイプラインを形成する後続の動作の一部になります。

インテル^® FPGA SDK for OpenCL™ ホストはSingle Work-Itemとしてカーネルを実行できます。これはNDRangeサイズが（1,1,1）のカーネルを起動するのと同じです。

OpenCL仕様バージョン1.0では、この動作モードがタスク・パラレル・プログラミングとして記述されています。 タスクとは、Single Work-Itemを含む1つのワークグループで実行されるカーネルのことです。

一般に、ホストはMultiple Work-Itemを並行して起動します。しかし、このデータ並列プログラミング・モデルは、並列Work-Item間で細かいデータを共有する必要がある場合には適していません。このような場合、カーネルをSingle Work-Itemとして表現することで、スループットを最大化することができます。 NDRangeカーネルとは異なり、Single Work-Itemカーネルは、Cプログラミングに似た自然なシーケンシャル・モデルに従います。特に、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 }

図 6. ループ解析レポート

図 7. Single Work-Itemカーネルのシステムビュー

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

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

図 9. Single Work-Itemの実行

NDRangeで書かれた同じ累算例を検討してください。

kernel void accum_ndr (global int* a, 
                       global int* c, 
                       int size) {
   int k = get_global_id(0);

   int sum[1024];
   for (int i = 0; i < size; ++i) {
     int j = k * size + i;
     sum[k] += a[j];
   }
   c[k] = sum[k];
}

図 10. ループ解析レポート

図 11. System View of the NDRange Kernel

パラレル・スレッドは、クロックサイクルごとに1つのスレッドでパイプライン方式で起動されます。この場合、ループ・パイプラインでパイプラインの並列処理とループ反復間の状態情報の通信が可能になります。ループの依存関係は、1クロックサイクルで解決されないことがあります。

以下の図は、シングル・スレッド・ホスト・アプリケーションが、カーネル実行間で並列の独立したデータパスをどのように処理するかを示しています。

シングル・スレッドのホスト・アプリケーションでは、 OpenCL™ホスト関数呼び出しの周りに外部同期機構を構築する必要があります。スレッドセーフな実行環境でマルチ・スレッド・ホスト・アプリケーションを使用すると、ホストコードを簡素化できます。さらに、ホスト内の複数のデータセットを同時に処理することで、カーネルの実行を高速化することができます。

次の図は、マルチ・スレッド・ホスト・アプリケーションがカーネル実行間で並列の独立したデータ・パスをどのように処理するかを示しています。

Report Summary は、デザイン内の各カーネル の要約を含んだデザインのコンパイル結果のおおまかな概要と、デザイン内の各カーネル が使用する見積られたリソースの要約を提供します。

Report Summary は、Info、Kernel Summary 、Estimated Resource Usage、および Compile Warnings の 4 つ のセクションに分かれています。

 1  // ND-Range kernel with unrolled loops
 2  __attribute((reqd_work_group_size(1024,1,1)))
 3  kernel void t (global int * out, int N) {
 4    int i = get_global_id(0);
 5    int j = 1;
 6    for (int k = 0; k < 4; k++) {
 7      #pragma unroll
 8      for (int n = 0; n < 4; n++) {
 9        j += out[k+n];
10      }
11    }
12    out[i] = j;
13
14    int m = 0;
15    #pragma unroll 1
16    for (int k = 0; k < N; k++) {
17      m += out[k/3];
18    }
19    #pragma unroll
20    for (int k = 0; k < 6; k++) {
21      m += out[k];
22    }
23    #pragma unroll 2
24    for (int k = 0; k < 6; k++) {
25      m += out[k];
26    }
27    out[2] = m;
28  }

このデザイン例のループ分析レポートは、コードで定義されているさまざまな種類のループのアンロール戦略を強調しています。

図 15. 4つのループを持つOpenCLデザイン例のループ解析レポート

Intel^® FPGA SDK for OpenCL™オフライン・コンパイラーはソースコードに基づいて以下のループ・アンローリング戦略を実行します。

• 最初のループ（ライン6）を完全に展開します。
• #pragma unrollの仕様のため、最初のループ内（ライン8）で内側のループを完全に展開します。
• #pragma unroll 1仕様のため、2番目の外部ループBlock2（ライン16）をアンロールしません。
• #pragma unrollの仕様のため、3番目の外部ループ（ライン20）を完全に展開します。
• #pragma unroll 2の仕様のため、4番目の外部ループBlock4（ライン24）を2回展開します。

kernel void big_lmem_4r_4w_nosplit (global int* restrict in, 
                                    global int* restrict out) {  
    local int lmem[4][1024]; 
 
    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 < 4; i++) {          
         lmem[i][(li*i) % ls] = res;    
         res >>= 1;  }    

    // Global memory barrier
    barrier(CLK_GLOBAL_MEM_FENCE); 

    res = 0;  
    #pragma unroll   
    for (int i = 0; i < 4; i++) {    
        res ^= lmem[i][((ls-li)*i) % ls];  }     
    out[gi] = res;
}

この例のシステム・ビューア・レポートは、ストール可能なロードおよびストアを強調表示します。

ロードとストアの動作の間の最初のバンクでは、アービトレーションが高い2つのメモリーバンクしか作成されないことに注意してください。次に、次のコード例に示すように、バンク・インデックスを第2次元に切り替えます。

kernel void big_lmem_4r_4w_nosplit (global int* restrict in, 
                                    global int* restrict out) {  
  local int lmem[1024][4];  

  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 < 4; i++) {          
    lmem[(li*i) % ls][i] = res;    
    res >>= 1; 
  }    
    
  // Global memory barrier
  barrier(CLK_GLOBAL_MEM_FENCE); 
    
  res = 0;  
  #pragma unroll   
  for (int i = 0; i < 4; i++) {    
    res ^= lmem[((ls-li)*i) % ls][i];  
  }     
  out[gi] = res;
}

カーネル・メモリー・ビューアでは、4つのメモリーバンクが別々のロード・ストア・ユニットで作成されていることがわかります。すべてのロードストア命令はストールフリーです。

origカーネルはネストされたループを4の深さまでがあります。ネストされたループは、次のレポートに示すように、変数が保持されているためにエリアを消費する余分なブロック（ブロック2,3,4,6,7,8）を作成しました。

ループの合体により、 lc_testのレイテンシーが短縮されていることがわかります。 origカーネルのブロック5とlc_testカーネルのブロック12は、最も内側のループです。

4つのループを含むOpenCLカーネルの例：

 1  // ND-Range kernel with unrolled loops
 2  __attribute ((reqd_work_group_size(1024,1,1)))
 3  kernel void t (global int * out, int N) {
 4    int i = get_global_id(0);
 5    int j = 1;
 6    for (int k = 0; k < 4; k++) {
 7      #pragma unroll
 8      for (int n = 0; n < 4; n++) {
 9        j += out[k+n];
10      }
11    }
12    out[i] = j;
13
14    int m = 0;
15    #pragma unroll 1
16    for (int k = 0; k < N; k++) {
17      m += out[k/3];
18    }
19    #pragma unroll
20    for (int k = 0; k < 6; k++) {
21      m += out[k];
22    }
23    #pragma unroll 2
24    for (int k = 0; k < 6; k++) {
25      m += out[k];
26    }
27    out[2] = m;
28  }

以下のエリアレポートには、カーネルシステム、ボード・インターフェイス、およびグローバル・インターコネクトのエリア使用量がリストされています。これらの要素はシステムレベルのIPで、デザインがターゲットとするカスタム・プラットフォームまたはリファレンス・プラットフォームに依存します。カーネルtはカーネルシステムの階層内にあり、ソースコードが始まる場所です。このレポートは、カーネルtの下でソースコード内で宣言されたすべての変数を指定し、残りのエリア情報を行番号でソートします。

図 24. 例示的なエリアレポートのソースビュー

この例では、コードラインj + = out [k + n]のコードライン（9行目）では、 Intel^® FPGA SDK for OpenCL™オフライン・コンパイラー追加を実行してグローバルメモリーからデータをロードするのに必要なエリアに基づいて推定エリア使用量を計算します。コードラインout [i] = j （12行目）の場合、オフライン・コンパイラーはポインター値を計算するために必要なエリアに基づいて推定エリア使用量を計算し、グローバルメモリーに格納し直します。

4つのループを含むOpenCLカーネルの例：

 1  // ND-Range kernel with unrolled loops
 2  __attribute((reqd_work_group_size(1024,1,1)))
 3  kernel void t (global int * out, int N) {
 4    int i = get_global_id(0);
 5    int j = 1;
 6    for (int k = 0; k < 4; k++) {
 7      #pragma unroll
 8      for (int n = 0; n < 4; n++) {
 9        j += out[k+n];
10      }
11    }
12    out[i] = j;
13
14    int m = 0;
15    #pragma unroll 1
16    for (int k = 0; k < N; k++) {
17      m += out[k/3];
18    }
19    #pragma unroll
20    for (int k = 0; k < 6; k++) {
21      m += out[k];
22    }
23    #pragma unroll 2
24    for (int k = 0; k < 6; k++) {
25      m += out[k];
26    }
27    out[2] = m;
28  }

図 25. 地域レポートの例

システムビューでは、カーネルは論理ブロックに分割されています。ブロックに関連付けられているコードラインのエリア使用情報を表示するには、そのブロックのレポートエントリーを展開するだけです。この例では、コードラインout [i] = j （12行目）のエリア情報がBlock1の下で使用可能です。システムビューでのライン12の推定エリア使用量は、ソースビューでの見積もりと同じです。

次の方法でシステムビューアと対話できます。

• マウスホイールを使用して、システムビューア内で拡大または縮小します。
• 赤い論理ブロックに関連付けられているデザインの部分を確認します。たとえば、High Initiation Interval (II）値を持つパイプライン・ループを持つロジックブロックは、High II値がデザインのスループットに影響を及ぼす可能性があるため、赤で強調表示されます。
  
• ブロック内の任意のノードにカーソルを合わせると、そのノードの情報がツールチップと詳細ペインに表示されます。
  

• 非表示にする接続のタイプをオフにして、システムビューアに含める接続のタイプを選択します。デフォルトでは、 ControlとMemoryの両方がシステムビューアでチェックされます。 Controlとは、ブロックとループの間の接続を指します。Memoryとは、グローバルメモリーまたはローカルメモリーとの接続を指します。デザインに読み出しチャンネルまたは書き込みチャンネルとの接続がある場合、システムビューアにChannelsオプションもあります。

多くのアルゴリズムでは、データ移動がボトルネックになることがよくあります。High Level Design レポート (report.html) 内の Kernel memory viewer は、 Intel^® FPGA SDK for OpenCL™オフライン・コンパイラーが カーネルのメモリーシステム上のデータ接続をどのように解釈するかを示しています。 kernel memory viewer を使用すると、 カーネルデザインのデータ移動のボトルネックの特定に役立ちます。

更に、メモリーアクセスの一部のパターンは LSU ( ロードストアー・ユニット ) で望まない調停を引き起こし、 カーネルのスループット性能に影響を与える可能性があります。 Kernel memory viewer を使用すると、LSU で望ましくない調停が発生する可能性のある位置を見つけることができます。

行列の正方形AxAを実行するOpenCLのデザイン例：

// performs matrix square A*A
// A is a square len*len matrix

kernel void matrix_square (global float* restrict A, 
                           unsigned len, 
                           global float* restrict out) {
  for(unsigned oi = 0; oi < len*len; oi++) {
    float sum = 0;
    int row = oi % len;
    for (int col = 0; col < len; col++) {
      unsigned i = (row * len) + col;  // 0, 1, 2, 3, 4,...
      unsigned j = (col * len) + row;  // 0, 3, 6, 9, 1,...
      sum += A[i] * A[j];
    }
    out[oi] = sum;
  }
}

カーネルmatrix_squareのエリアレポートのシステムビューは、 ブロック3のフリップフロップ（FF）とRAMの推定使用量が高いことを示しています。システムビューアのBlock3をさらに調べると、Block3のレイテンシー値も高いことがわかります。

図 26. 最適化されていないカーネルmatrix_squareのエリアレポートとシステムビューアの結果

これらのパフォーマンスのボトルネックの原因は、システムがループ内からグローバルメモリーからデータをロードしているためです。したがって、次の変更されたコードに示すように、データをローカルメモリーにプリロードすることが、最初に行う最適化のステップです。

kernel void matrix_square_v1 (global float* restrict A, 
                              unsigned len, 
                              global float* restrict out)
{
  // 1. preload the data into local memory
  //    - suppose we know the max size is 4X4

  local int cache_a[16];for(unsigned k = 0; k < len*len; k++) {  cache_a[k] = A[k];}

  for(unsigned oi = 0; oi < len*len; oi++) {
    float sum = 0;
    int row = oi % len;
    for(int col = 0; col < len; col++) {
      unsigned i = (row * len) + col;  // 0, 1, 2, 3, 4,...
      unsigned j = (col * len) + row;  // 0, 3, 6, 9, 1,...
      sum += cache_a[i] * cache_a[j];
    }
    out[oi] = sum;
  }
}

図 27. 修正されたカーネルmatrix_square_v1のエリアレポートとシステムビューアの結果

エリアレポートとシステムビューアの結果に示されているように、ローカルメモリーにデータをプリロードすると、RAMの使用量が3分の1に減少し、レイテンシー値が255から97に低下します。

matrix_square_v1のエリアレポートをさらに調べると、以下のエリアレポートの行30であるコードラインint row = oi％lenは、法計算のために異常に大きなエリアが使用されます。

図 28. 修正カーネルmatrix_square_v1のエリアレポートのソースビュー

モジュラス計算を削除して列カウンターに置き換えると、修正されたカーネルmatrix_square_v2に示すように、適応ルックアップテーブル（ALUT）およびFF使用量を50％削減できます。

kernel void matrix_square_v2 (global float* restrict A, 
                              unsigned len, 
                              global float* restrict out)
{
  // 1. preload the data into local memory
  //    - suppose we know the max size is 4X4
  // 2. remove the modulus computation

  local int cache_a[16];
  for (unsigned k = 0; k < len*len; k++) {
    cache_a[k] = A[k];
  }

  unsigned row = 0;
  unsigned ci = 0;
  for (unsigned oi = 0; oi < len*len; oi++) {
    float sum = 0;

    // keep a column counter to know when to increment rowif (ci == len) {  ci = 0;  row += 1;}ci += 1;

    for (int col = 0; col < len; col++) {
      unsigned i = (row * len) + col;   // 0, 1, 2, 3, 4,...
      unsigned j = (col * len) + row;   // 0, 3, 6, 9, 1,...
      sum += cache_a[i] * cache_a[j];
    }
    out[oi] = sum;
  }
}

図 29. 修正されたカーネルmatrix_square_v1とmatrix_square_v2間のエリア使用の比較

matrix_square_v2のエリアレポートをさらに調べると、インデックスiとjの計算（つまり、符号なしi =（row * len）+ colおよび符号なしj =（col * len）+ row）ではALUTとFFの使用量の見積もりが非常に異なることがわかります。さらに、エリアレポートは、これらの2つの計算がデジタル信号処理（DSP）ブロックを使用していることも示しています。

図 30. 異なるインデックス計算の異なるエリアの使用を示す修正カーネルmatrix_square_v2のエリアレポート

インデックス計算のためにDSPとRAMブロックの使用を最適化する方法は、乗算計算を削除して、下記の修正されたカーネルmatrix_square_v3に示すように、加算を追跡するだけです。

kernel void matrix_square_v3 (global float* restrict A, 
                              unsigned len, 
                              global float* restrict out) {
  // 1. preload the data into local memory
  //    - suppose we know the max size is 4X4
  // 2. remove the modulus computation
  // 3. remove DSP and RAM blocks for index calculation helps reduce the latency

  local int cache_a[16];
  for (unsigned k = 0; k < len*len; k++) {
    cache_a[k] = A[k];
  }

  unsigned row_i = 0;
  unsigned row_j = 0;
  unsigned ci = 0;
  for (unsigned oi = 0; oi < len*len; oi++) {
    float sum = 0;
    unsigned i, j;

    // keep a column counter to know when to increment row
    if (ci == len) {
      ci = 0;
      row_i += len;
      row_j += 1;
    }
    ci += 1;

    i = row_i;  // initialize i and jj = row_j;for (int col = 0; col < len; col++) {i += 1;       // 0, 1, 2, 3, 0,...j += len;     // 0, 3, 6, 9, 1,...
      sum += cache_a[i] * cache_a[j];
    }
    out[oi] = sum;
  }
}

乗算ステップを削除することで、下記のエリアレポートに示すように、DSP使用率を50％削減できます。さらに、この修正はレイテンシーを短縮するのに役立ちます。

図 31. 修正されたカーネルmatrix_square_v2とmatrix_square_v3間のエリア使用の比較

レイテンシーをさらに削減するために、修正されたカーネルmatrix_square_v3のループ分析レポートを確認することができます。以下に示すように、解析ペインと詳細ペインでは、 sum + = cache_a [i] * cache_a [j]のループに依存する依存関係があるため、Block27にIIボトルネックが発生しています。

図 32. 修正されたカーネルのループ解析matrix_square_v3

ループで運ばれる依存関係を解決するには、修正されたカーネルmatrix_square_v4で強調表示されているコードに示すように、計算の乗算部分と加算部分を分けることができます。

kernel void matrix_square_v4 (global float* restrict A, 
                              unsigned len, 
                              global float* restrict out)
{
  // 1. preload the data into local memory
  //    - suppose we know the max size is 4X4
  // 2. remove the modulus computation
  // 3. remove DSP and RAM blocks for index calculation helps reduce the latency
  // 4. remove loop-carried dependency 'sum' to improve throughput by trading off area

  local int cache_a[16];
  for (unsigned k = 0; k < len*len; k++) {
    cache_a[k] = A[k];
  }

  unsigned row_i = 0;
  unsigned row_j = 0;
  unsigned ci = 0;
  for (unsigned oi = 0; oi < len*len; oi++) {
    float sum = 0;
    unsigned i, j;

    float prod[4];   // make register
    #pragma unroll
    for (unsigned k = 0; k < 4; k++) {
      prod[k] = 0;
    }

    // keep a column counter to know when to increment row
    if (ci == len) {
      ci = 0;
      row_i += len;
      row_j += 1;
    }
    ci += 1;

    i = row_i;  // initialize i and j
    j = row_j;
    for (int col = 0; col < len; col++) {
      i += 1;       // 0, 1, 2, 3, 0,...
      j += len;     // 0, 3, 6, 9, 1,...
      prod[col] = cache_a[i] * cache_a[j];
    }

    sum = prod[0];#pragma unrollfor (unsigned k = 1; k < 4; k++) {  sum += prod[k];}

    out[oi] = sum;
  }
}

以下のエリアレポートおよびシステムビューアの結果に示されているように、計算ステップを分割することで、エリア使用量の増加を犠牲にしてより高いスループットを達成できます。この変更により、ループのII値が1に減少し、レイテンシーが30から24に減少します。

図 33. 修正されたカーネルmatrix_square_v4のエリアレポートとシステムビューアの結果

組み込みのWork-Item関数の詳細については、OpenCL Specification version 1.0の6.11.1: Work-Item Functionsのいセクションを参照してください。

単一のWork-Itemカーネルの場合、オフライン・コンパイラーはカーネル内のすべてのループをパイプライン化して、複数のループ反復を同時に実行できるようにします。コンパイラーがループの一部を効果的にパイプライン化できない場合、またはループをパイプライン化できない場合、カーネルのパフォーマンスが低下する可能性があります。

オフライン・コンパイラーは、NDRangeカーネルでループをパイプライン化できません。ただし、これらのループは複数のWork-Itemを同時に受け入れることができます。カーネルには複数のループがあり、それぞれにネストされたループがあります。外側のループごとにネストされたループの反復の総数を表にすると、カーネルのスループットは、通常、テーブル化した最大の反復の合計値だけ減少します。

効率的にNDRangeカーネルを実行するには、通常、多数のスレッドが必要です。

GPUとは異なり、FPGAはアプリケーションに最適なカスタムLSUを構築できます。結果として、アプリケーションに理想的なLSUタイプを選択するOpenCLコードを書く能力が、デザインのパフォーマンスを大幅に向上させるのに役立つかもしれません。

デザインのHTMLエリアレポートを見直すと、システムレベルのグローバル・インターコネクトエントリーの値は、グローバル・メモリー・インターコネクトのサイズを表します。

図 35. OpenCLデザインにおけるグローバル・メモリー・インターコネクトのサイズを示すHTMLエリアレポート

HTMLレポートでは、メモリー・システム・ビューアは、グローバルメモリー相互接続を負荷（LD）、ストア（ST）、および接続（灰色線）として示しています。

図 36. OpenCLデザインにおけるグローバルメモリー相互接続のシステムビューア結果

Intel^® FPGA SDK for OpenCL Offline Compilerは、デザインのメモリー・アクセス・パターンに基づいて、OpenCLシステム用の適切なタイプのLSUを選択します。例示的なLSUタイプには、連続アクセス（または連続アクセス）およびバースト・インターリーブ・アクセスが含まれています。 図 79そして図 78連続したメモリーアクセスとバースト・インターリーブされたメモリーアクセスとの間のアクセスパターンの違いをそれぞれ示しています。

カーネルコードでチャネルを宣言するときは、宣言の前にキーワードchannelを付けてください 。

例：channel long16 myCh __attribute__((depth(16)));

HTMLレポートでは、エリアレポートによってチャネルエリアがソースコードの宣言行にマップされます。チャネルおよびチャネルアレイは、その幅および深さとともに報告されます。

実装されたチャネル深度は、チャネル宣言で指定した深度と異なる場合があります。 OpenCLオフライン・コンパイラー用インテルFPGA SDKは、シフトレジスターまたはRAMブロックにチャネルを実装できます。オフライン・コンパイラーは、チャネル深度に基づいてチャネル実装のタイプを決定します。

Intel^® FPGA SDK for OpenCL™オフライン・コンパイラーさまざまな種類のロードストアユニット（LSU）を生成します。 LSUの種類によっては、コンパイラーがメモリー・アクセス・パターンやその他のメモリー属性に応じてLSUの動作やプロパティを変更することがあります。

ロードストアのユニットタイプまたは修飾子を明示的に選択することはできませんが、コード内のメモリー・アクセス・パターン、使用可能なメモリーのタイプ、およびメモリーアクセスがローカルメモリーかグローバルメモリーかを変更することによって、コンパイラーがインスタンス化するLSUのタイプに影響を与えることができます 。

一般に、最初に単一のコンピューティング・ユニットをターゲットとするカーネルを最適化する必要があります。このコンピューティング・ユニットを最適化した後、ハードウェアをスケーリングしてFPGAの残りの部分を満たすようにパフォーマンスを上げてください。カーネルのハードウェア・フットプリントは、ハードウェアのコンパイルに要する時間と相関します。したがって、より小さなフットプリント（つまり、単一の計算単位）で実行できる最適化が増えるほど、一定の時間内に実行できるハードウェアの数が増えます。

データ処理とメモリーアクセスの最適化に加えて、カーネルの作成時には、必要に応じて次のデザイン方法を実装することを検討してください。

時には、FPGAからグローバルメモリーへの帯域幅は、カーネル間のデータ転送効率を制限します。理論上の最大FPGA対グローバルメモリー帯域幅は、対象のカスタム・プラットフォームおよびボードで使用可能なグローバルメモリーバンクの数によって異なります。ボードの理論上の最大帯域幅を決定するには、ボードベンダのマニュアルを参照してください。

実際には、カーネルは使用可能な最大グローバルメモリー帯域幅の100％使用を達成していません。使用率は、アルゴリズムのアクセスパターンによって異なります。

グローバルメモリー帯域幅がOpenCLカーネルのパフォーマンス制約条件である場合、まずアルゴリズムを複数の小さなカーネルに分解してみてください。次に、下の図に示すように、SDKカーネル間のデータ転送用のチャネルまたはOpenCLパイプを実装してグローバル・メモリー・アクセスの一部を削除します。

チャネルの使用方法の詳細については インテル^® FPGA SDK for OpenCL™ プログラミング・ガイドの インテル^® FPGA SDK for OpenCL™ チャネル拡張の実装のセクションを参照してください。

パイプの使用方法の詳細については、 インテル^® FPGA SDK for OpenCL™ プログラミング・ガイドのOpenCLパイプの実装のセクションを参照してください。

次のコードの例を検討してみましょう。

__kernel void
consumer (__global uint*restrict dst) {
  for (int i = 0; i < 5; i++) {
    dst[2*i] = read_channel_intel(c0);
    dst[2*i+2] = read_channel_intel(c1);
  }
}

__kernel void
consumer (__global uint*restrict dst,
  read_only pipe uint 
    __attribute__((blocking)) c0,
  read_only pipe uint
    __attribute__((blocking)) c1)
{
  for (int i = 0; i < 5; i++) {
    read_pipe (c0, &dst[2*i]);
    read_pipe (c1, &dst[2*i+2]);
  }
}

左側のコード例は、2つの読み出しチャネル呼び出しを行います。右側のコード例は、2つの読み出しパイプ呼び出しを行います。ほとんどの場合、カーネルはこれらのチャネル呼び出しまたはパイプ呼び出しを並列に実行します。ただし、チャネルおよびパイプ・コールの実行が順不同で行われる可能性があります。アウト・オブ・シーケンス実行は、C1からの読み出し動作が発生し、C0からの読み出し動作の前に完了することができることを意味します。

コンパイル時に、オフライン・コンパイラーは、相互作用するチャネルまたはパイプ間のスケジューリングの不一致を計算します。これらの不一致は、読み出しと書き込みの動作の不均衡を引き起こす可能性があります。オフライン・コンパイラーは、不均衡を修正するためにバッファー推論最適化を自動的に実行します。

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

__kernel void producer (
  __global const uint * restrict src,
  const uint iterations)
{
  for(int i = 0; i < iteration; i++)
  {
    write_channel_intel(c0,src[2*i]);
    write_channel_intel(c1,src[2*i+1]);
  }
}

__kernel void consumer (
  __global uint * restrict dst,
  const uint iterations)
{
  for(int i = 0; i < iterations; i++)
  {
    dst[2*i] = read_channel_intel(c0);
    dst[2*i+1] = read_channel_intel(c1);
  }
}

__kernel void producer (
  __global const uint * restrict src,
  const uint iterations,
  write_only pipe uint
    __attribute__((blocking)) c0,
  write_only pipe uint
    __attribute__((blocking)) c1)
{
  for(int i = 0; i < iteration; i++)
  {
    write_pipe(c0,&src[2*i]);
    write_pipe(c1,&src[2*i+1]);
  }
}

__kernel void consumer (
  __global uint * restrict dst,
  const uint iterations,
  read_only pipe uint
    __attribute__((blocking)) c0,
  read_only pipe uint
    __attribute__((blocking)) c1)
{
  for(int i = 0; i < iterations; i++)
  {
    read_pipe(c0,&dst[2*i]);
    read_pipe(c1,&dst[2*i+1]);
  }
}

オフライン・コンパイラーは、カーネル間のチャネルまたはパイプがサイクルを形成できない場合、バッファー推論の最適化を実行します。カーネル間のcyclは、カーネルから、書き込みチャネルまたは書き込みパイプ呼び出しを経由して元のカーネルに戻るパスです。この例では、カーネルproducerのライトチャネルまたはライトパイプコールが10サイクル離れてスケジュールされ、リードチャネルまたはリードパイプコールが15サイクル離れてスケジュールされているとします。C1への読み出し動作が発生する前に、5つの余分な書き込み動作が発生するかもしれないのであり、読み出しが一時的に不一致が存在するとC1に書き込み動作を。この不均衡を修正するために、オフライン・コンパイラーは、ストールを避けるためにc1に5サイクルのバッファーサイズを割り当てます。追加のバッファー容量は、 producerカーネルのc1書き込み動作とconsumerカーネルのc1読み出し動作を切り離します。

• マルチ・スレッド・カーネルよりもシングル・スレッド・カーネルを使用してください。
• デザインモデルがフィード・フォワード・データパス、例えばバックツーバックループまたは離散処理ステップでどのように表現できるかを検討します。デザインをチャネルで接続された複数のカーネルに分割する必要があるかどうかを判断します。
• カーネルの同じポイントでデータ全体が使用されている場合のみ、チャネル上のデータを集約します。
• カーネルあたりのチャネル数を妥当なものに保つようにしてください。
• データを待っているループ構造を使用している場合、非ブロッキング・チャネルまたはパイプを使用しないでください。非ブロッキング・チャネルは、ブロッキング・チャネルより多くのリソースを消費します。

各Work-Itemが配列内の4つの要素の累積を計算する責任がある並列アプリケーション用のOpenCLコードを検討してください。

__kernel void example ( __global const int * restrict x,
                        __global int * restrict sum ) {
   int accum = 0;

   for (size_t i = 0; i < 4; i++) {
      accum += x[i + get_global_id(0) * 4];
   }

   sum[get_global_id(0)] = accum;
}

このカーネルでは、以下の3つの主要な動作が行われています。

• 入力xからのロード動作
• 累積
• オペレーションを出力sumに格納する

オフライン・コンパイラーは、OpenCLカーネルコードのデータフローセマンティクスに従って、これらの動作をパイプラインに配置します。たとえば、オフライン・コンパイラーは、ループ終了条件に応じて、パイプラインの最後からパイプラインの先頭に結果を転送することによってループを実装します。

OpenCLカーネルは、クロックサイクルごとに各Work-Itemの1回のループ反復を実行します。十分なハードウェア・リソースがあれば、ループを展開することでカーネルのパフォーマンスを向上させることができ、カーネルが実行する反復回数が減ります。ループをアンロールするには、次のコード例に示すように、 #pragma unrollディレクティブをメインループに追加します。ループをアンロールすると、オフライン・コンパイラーが作成するコンピューティング・ユニットの構造が大幅に変更されることに注意してください。

__kernel void example ( __global const int * restrict x,
                        __global int * restrict sum ) {
  int accum = 0;

  #pragma unroll
  for (size_t i = 0; i < 4; i++) {
    accum += x[i + get_global_id(0) * 4];
  }

  sum[get_global_id(0)] = accum;
}

この例では、 #pragma unrollディレクティブにより、オフライン・コンパイラーはループの4つの繰り返しを完全に展開します。アンロールを達成するために、オフライン・コンパイラーは、加算演算の数を3倍にし、4倍のデータをロードすることによってパイプラインを拡張します。ループが除去されると、コンピューティング・ユニットはフィードフォワード構造を前提します。結果として、コンピューティング・ユニットは、初期ロード動作および加算の完了後にクロックサイクルごとにsum要素を記憶することができる。オフライン・コンパイラーは、4つのロード動作を統合することによってこのカーネルをさらに最適化し、コンピューティング・ユニットが1つのロード動作で結果を計算するために必要なすべての入力データをロードできるようにします。

ループをアンロールし、グローバルメモリーからロード動作を統合することにより、ハードウェアによるカーネルのインプリメンテーションでクロックサイクルごとに多くの動作を実行できます。一般に、OpenCLカーネルのパフォーマンスを向上させるために使用する方法は、次の結果を達成する必要があります。

• 並列動作の数を増やす
• 実装のメモリー帯域幅を増やす
• カーネルがハードウェアで実行できるクロックサイクルあたりの動作数を増やす

オフライン・コンパイラーは、以下の状況で完全にループを展開することができない場合があります。

• 非常に多数の反復を伴うデータ依存ループの完全なアンローリングを指定します。したがって、カーネルのハードウェア実装がFPGAに適合しない可能性があります。
• 完全な展開を指定し、ループの境界は定数ではありません。
• ループは、複雑な制御フローで構成されます（たとえば、コンパイル時に不明な複雑な配列インデックスまたは終了条件を含むループ）。

上記の最後の2つのケースでは、オフライン・コンパイラーは次の警告を発行します。

ループのフルアンロールが要求されますが、ループの境界を特定できません。ループはアンロールされません。

これらの状況でループのアンロールをイネーブルするには、 #pragma unroll <N> ディレクティブ(<N>はアンロール係数である）を指定します。アンロール係数は、オフライン・コンパイラーがアンロールする回数を制限します。たとえば、カーネル内のループが展開されないようにするには、そのループに#pragma unroll 1ディレクティブを追加します。

適切に構成されたループを構築するためのヒントについては、Good Design Practices for Single Work-Item Kernelを参照してください。

OpenCL^®標準は固定小数点表現をサポートしていません。整数データ型を使用して固定小数点表現を実装する必要があります。ハードウェア開発者は、通常、固定小数点データ表現を使用してハードウェアを節約し、計算を実行するために必要なデータ解決のみを保持します。OpenCL標準ではこれらのデータ解像度のみがサポートされているため、8,16,32、または64ビットのスカラデータ型を使用する必要があります。ただし、ハードウェアのコンパイルツールでハードウェア・リソースを節約するための最適化を実行できるように、ソースコードに適切なマスキング動作を組み込むことができます。

たとえば、アルゴリズムで17ビットのデータの固定小数点表現を使用する場合、値を格納するために32ビットのデータ型を使用する必要があります。その後、 Intel^® FPGA SDK for OpenCL™オフライン・コンパイラー 2つの17ビット固定小数点値を一緒に追加するには、オフライン・コンパイラーが余分な上位15ビットの加算を処理するために余分なハードウェアを作成する必要があります。この追加のハードウェアが存在しないようにするには、スタティックビットマスクを使用して、ハードウェアのコンパイル中に不要なビットを無視するようにハードウェアのコンパイルツールに指示します。以下のコードはこのマスキング動作を実装しています。

__kernel fixed_point_add (__global const unsigned int * restrict a,
                          __global const unsigned int * restrict b,
                          __global unsigned int * restrict result)
{
	   size_t gid = get_global_id(0);

   	unsigned int temp;
   	temp = 0x3_FFFF & ((0x1_FFFF & a[gid]) + ((0x1_FFFF & b[gid]));

   	result[gid] = temp & 0x3_FFFF;
}

このコード例では、入力aとbの上位15ビットがマスクされて一緒に加算されます。 2つの17ビット値を加算した結果は18ビットの分解能を超えることができないため、オフライン・コンパイラーは追加のマスクを適用して結果の上位14ビットをマスクします。最終的なハードウェアの実装は、完全な32ビットの追加ではなく、17ビットの追加です。この例のロジックの節約は、FPGAで使用可能なハードウェア・リソースの数に比べて比較的少ないです。しかし、これらの小さな節約は、頻繁に適用されると、FPGA全体でより大きなハードウェア節約に蓄積されます。

アラインメントされたメモリー割り当てを設定するには、ホストプログラムに次のソースコードを追加します。

• Windowsの場合

  #define AOCL_ALIGNMENT 64
  #include <malloc.h>
  void *ptr = _aligned_malloc (size, AOCL_ALIGNMENT);

  整列メモリーブロックを解放するには、関数呼び出し_aligned_free(ptr);を含めます。

• Linuxの場合

  #define AOCL_ALIGNMENT 64
  #include <stdlib.h>
  void *ptr = NULL;
  posix_memalign (&ptr, AOCL_ALIGNMENT, size);

  整列メモリーブロックを解放するには、関数呼び出しfree(ptr);を含めます。

typedef struct {
	char r,g,b,alpha;
} Pixel_s;

typedef union {
     Pixel_s p;
     int not_used;
} Pixel;

typedef struct {char r、g、b、alpha; } __attribute __（（aligned（4）））ピクセル;

オフライン・コンパイラーは、次のすべての基準を満たすために構造体のアラインメントを必要とするISO C標準に準拠しています。

• アラインメントは、すべての構造体メンバーのアラインメント間の最小公倍数の整数倍でなければなりません。
• このマクロの値は、2の累乗である必要があります。

カーネルコードにaligned（ N ）属性を含めることで、 structの整列を設定することができます。整列属性なしで、オフライン・コンパイラーは、 structのサイズに基づいて、 structのアレイ内の各structのアラインメントを決定します。次の例を検討してください。

__kernel void test (struct mystruct* A,
                    struct mystruct* B)
{
    A[get_global_id(0)] = B[get_global_id(0)];
}

mystructのサイズが101バイトの場合、各ロードまたはストアのアクセスは1バイトで整列されます。 mystructのサイズが128バイトの場合、各ロードまたはストアのアクセスは128バイトに整列され、最も効率的なハードウェアが生成されます。

構造体のフィールドは、structの中に整列されていない場合、オフラインのコンパイラーはそれらを整列するためにパディングを挿入します。 structフィールド間にパディングを挿入すると、次のようにハードウェアの効率に影響します。

• 構造体のサイズを増やす
• アラインメントに影響する可能性がある

オフライン・コンパイラーがパディングを挿入しないようにするには、カーネルコードにpacked属性を含めます。前述のISO C標準は、パックされたstructまたはアンパックされたstructのアラインメントを決定するときに適用されます。次の例を検討してください。

struct mystruct1
{
    char a;
    int b;
};

mystruct1のサイズは8バイトです。したがって、 structは8バイトに整列され、カーネル内で効率的にアクセスされます。ここで別の例を検討してください。

struct mystruct2
{
    char a;
    int b;
    int c;
};

mystruct2のサイズは12バイトで、 structは4バイトに整列しています。structフィールドはパディングされており、structはアラインメントされていないため、カーネル内のアクセスは非効率的です。

以下にpacked属性を含むstructの例を示します。

struct __attribute__((packed)) mystruct3
{
    char a;
    int b;
    int c;
};

mystruct4のサイズは16バイトです。 mystruct4が整列され、 structフィールドの間にパディングがないため、このカーネルのアクセスはmystruct3のアクセスよりも効率的です 。

構造体にaligned（ N ）属性とpacked属性の両方を含めるには、次の例を考えてください。

struct __attribute__((packed)) __attribute__((aligned(16))) mystruct5
{
    char a;
    int b;
    int c;
};

mystruct5のサイズは9バイトです。 aligned（16）属性のため、 structは配列内の16バイト整列アドレスに格納されます。 mystruct5は16バイトで整列されており、パディングもないため、このカーネルでのアクセスは効率的です。

structの整列とaligned（ N ）およびpacked属性の詳細については、以下の文書を参照してください。

• OpenCL仕様バージョン1.2.1のセクション6.11.1
• インテル^® FPGA SDK for OpenCL™ プログラミング・ガイドのデータ構造の挿入のディセーブル
• Structアラインメントの指定のセクションの インテル^® FPGA SDK for OpenCL™ プログラミング・ガイド

次のコード例は、ベクトル要素を更新するシナリオを示しています。

__kernel void update (__global const float4 * restrict in,
                      __global const float4 * restrict out)
{
	  size_t gid = get_global_id(0);

	  out[gid].x = process(in[gid].x);
	  out[gid].y = process(in[gid].y);
	  out[gid].z = process(in[gid].z);
	  out[gid].w = 0; //Update w even if that variable is not required.
}

restrictキーワードは、オフライン・コンパイラーに、ポインターが別のポインターのエイリアスではないことを通知します。たとえば、カーネルに大域メモリーAとBが重複しない2つのポインターがある場合、カーネルを次のように宣言します。

__kernel void myKernel (__global int * restrict A,
                        __global int * restrict B)

以下の機能は高価です。

• 整数除算とモジュロ（剰余）演算子
• 加算、乗算、絶対値、および比較を除くほとんどの浮動小数点演算子　
  注: 浮動小数点演算の最適化の詳細については、「浮動小数点演算の最適化のセクションを参照してください。
• アトミック関数

対照的に、安価な機能はカーネルの性能にはほとんど影響を与えず、その実装は最小のハードウェアしか消費しません。

以下の機能は安価です。

• AND、NAND、OR、NOR、XOR、およびXNORなどの2進論理演算
• 1つの定数引数による論理演算
• 定数でシフト
• 整数の乗算と2の累乗である定数による除算

高価な関数がワークグループ内のすべてのWork-Itemに対して新しいデータを生成する場合、それをカーネルでコード化することは有益です。これに対して、以下のコード例は、NDRangeの各Work-Itemで実行される高価な浮動小数点演算（除算）のケースを示しています。

__kernel void myKernel (__global const float * restrict a,
                        __global float * restrict b,
                        const float c, const float d)
{
   size_t gid = get_global_id(0);
   
   //inefficient since each work-item must calculate c divided by d
   b[gid] = a[gid] * (c / d); 
}

この計算の結果は常に同じです。このような冗長かつハードウェアのリソース集約的な動作を回避するには、ホスト・アプリケーションで計算を実行し、使用するNDRange内のすべてのWork-Itemの引数として結果をカーネルに渡します。変更されたコードを以下に示します。

__kernel void myKernel (__global const float * restrict a,
                        __global float * restrict b,
                        const float c_divided_by_d)
{
   size_t gid = get_global_id(0);

   /*host calculates c divided by d once and passes it into  
   kernel to avoid redundant expensive calculations*/   
   b[gid] = a[gid] * c_divided_by_d;  						
}

Intel^® FPGA SDK for OpenCL™オフライン・コンパイラーは、NDRange全体でWork-Itemに依存しない動作を1つの動作に統合します。次に、すべてのWork-Itemにわたって結果を共有します。最初のコード例では、オフライン・コンパイラーは、 cによるdの除算がすべてのWork-Itemにわたって一定であるため、すべてのWork-Itemで共有される1つの除算ブロックを作成します。この最適化は冗長ハードウェアの量を最小限に抑えるのに役立ちますしかし、整数除算の実装には、相当量のハードウェア・リソースが必要です。したがって、除算演算をホスト・プロセッサーにオフロードし、結果を引数としてカーネルに渡してハードウェア・リソースを節約することが有益です。

例えば、次のコードは、そのようなget_global_id又はget_local_idなどのWork-ItemIDを含む分岐示します。

for (size_t i = 0; i < get_global_id(0); i++)
{
	  // statements
}

次のOpenCLカーネルプログラムを検討してください。

__kernel void add (__global int * a,
                   __global int * b, 
                   __global int * c)
{
    int gid = get_global_id(0);
    c[gid] = a[gid]+b[gid];
}

下の図に示すように、Profiler計測器は、カーネルプログラム用に生成されたパイプライン全体で、デイジーチェインでパフォーマンス・カウンターを接続します。ホストは、これらのカウンターによって収集されたデータを読み出します。たとえば、 PCI Express^® （ PCIe^® ）ベースのシステムでは、ホストはPCIe制御レジスターアクセス（CRA）または制御およびステータスレジスター（CSR）ポートを介してデータを読み出します。

Work-Itemの実行停止は、 インテル^® FPGA SDK for OpenCL™ パイプラインのさまざまな段階で発生する可能性があります。大量のメモリーアクセスまたはロードおよびストア動作を伴うアプリケーションは、メモリー転送の完了を可能にするために頻繁に停止することがあります。Profilerは、カーネル・パイプライン内の大部分の停止を引き起こすロードおよびストア動作またはチャネルアクセスを識別するのに役立ちます。

Intel FPGA Dynamic Profiler for OpenCLの使用方法については、 インテル^® FPGA SDK for OpenCL™ プログラミング・ガイドの「OpenCLカーネルのプロファイリングのセクションを参照してください。

• パフォーマンス・カウンターをカーネルに挿入するには、開発中にaocコマンドに-profile Intel^® FPGA SDK for OpenCL™オフライン・コンパイラーコマンドオプションを含めます。
• Profilerを使用せずにカーネルの_fmaxとパフォーマンスを定期的にチェックしてください。
• プロファイラーのオーバーヘッドを減らすには、ローカルフォルダーからホスト・アプリケーションを実行します。リモートまたはNASフォルダからホストを実行しないでください。
• カーネルの実行時間が20 msより長いことを確認してください。そうしないと、Profilerのオーバーヘッドが引き継がれます。
• すべてのロードおよびストア動作とチャネルがデータフローでどのように接続されているかを理解します。

たとえば、カーネルにイーサネットI/Oに対するチャネル読み出し呼び出しがあり、Profilerがストールを識別した場合、書込みチャネルは、イーサネットI/Oへのデータ書込みがカーネルの読み出し速度と同じ速度で行われていないことを意味します。

カーネル間チャネルの場合、チャネルの読み出し側と書き込み側の間に不均衡がある場合、または読み出しと書き込みのカーネルが同時に実行されていない場合にストールが発生します。

たとえば、読み出しが書き込みを実行するカーネルと同時に起動されないか、または読み出し操作が書き込み操作よりもはるかに遅い場合、Profilerは書き込みカーネル内のwrite_channel_intel呼び出しの停止を識別します。

Intel FPGA Dynamic Profiler for OpenCL GUIでのSource Codタブは、占有率（Occupancy％）カラムのツールチップがActivityの割合を指定する場合があります。Activityは、以下に説明するように、アクティビティーが述語に関連する点で占有とは異なります。

各LSUには、ivalid信号の他に述語信号があります。 ivalidシグナルは、上流のロジックが有効なデータをLSUに提供していることを示します。述語信号は、LSUが受信するデータに作用するべきであることを示します。Work-Itemまたはループ反復は、述語であってもメモリー命令を占有することができます。分岐ステートメントにループが含まれていない場合、オフライン・コンパイラーは分岐を変換して制御フローを最小限に抑え、より効率的なハードウェアを実現します。変換の一部として、メモリーおよびチャネル命令を述語化し、出力結果をマルチプレクサ論理によって選択する必要があります。

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

int addr = compute_address();
int x = 0;
if (some_rare_condition)
    x = src[addr];

オフライン・コンパイラーは、コードを次のように変更します。

int addr = compute_address();
int x = 0;
x = src[addr] if some_rare_condition;

この場合、 src []はクロックサイクルごとに有効なアドレスを受け取ります。パイプラインにSRC []自体は生成されないストールを想定すると、[] SRC用ivalid信号は、ほとんどの時間を高くなります。実際には、 src []は述語信号some_rare_conditionが真の場合にのみロードを実行します。したがって、この負荷動作では、占有率は高くなりますが、活動は低くなります。

ツールヒントで使用可能なアクティビティーの割合は、述語アクセスを考慮しないため、低いアクティビティーの割合に基づいて述語命令を識別できます。活動率が低いにもかかわらず、これらの指示には高い占有率があるかもしれません。

Intel FPGA Dynamic Profiler for OpenCL GUIでのSource Codeタブは、 Attributesカラムのツールヒントがキャッシュヒット率を指定することがあります。一部のグローバル負荷ユニットでは、 Intel^® FPGA SDK for OpenCL™オフライン・コンパイラーはプライベート・キャッシュをインスタンス化することがあります。この場合、オフライン・コンパイラーはこのキャッシュの有効性を測定するために追加のハードウェアカウンターを作成します。このプライベート・キャッシュの詳細は、HTMLエリアレポートで確認できます。

帯域幅の使用が非効率的である場合、また、アプリケーションの実行中に大量のデータ転送が必要な場合、メモリー命令が頻繁に停止します。非効率的なメモリーアクセスは、帯域幅の最適使用に至りません。そのような場合、カーネルのメモリーアクセスを分析して、改善が可能かどうかを確認してください。

チャネルへの読み出しアクセスと書き込みアクセスとの間に強い不均衡が存在する場合、チャネル命令はストールします。不均衡は、チャネル読み出しまたは異なる速度で動作する書き込みによって引き起こされる可能性があります。

たとえば、書き込みチャネル呼び出しのストール率が高いと判明した場合、読み出しチャネル呼び出しの占有率および活動度が低いかどうかを確認します。そうである場合、読み出しチャネル呼び出しを制御するカーネルの実行速度は、書き込みチャネル呼び出しを制御するカーネルにとって遅すぎるため、パフォーマンスのボトルネックにつながります。

メモリーまたはチャネルアクセスがパーセンテージの高いパイプラインストールを引き起こしている場合、メモリーまたはチャネルを指示するソースコードのラインは赤で強調表示されます。ストール率が20％以上になると、ストールの識別が高くなります。ストール率が高いほど、赤いハイライトが暗くなります。高失速率の値を簡単にトラバースするには、[ソースコード]タブの右下隅に右矢印と左矢印があります。

図 70. 高い失速率の識別

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

__kernel void proc (__global int * a, ...) {
  for (int i = 0; i < N; i++) {
    for (int j = 0; j < 1000; j++) {
      write_channel_intel (c0, data0);
    }
    for (int k = 0; k < 3; k++) {
      write_channel_intel (c1, data1);
    }
  }
}    

すべてのループがパイプライン化されていると仮定すると、トリップカウントが1000の最初の内部ループがクリティカルループです。トリップカウントが3の2番目の内部ループは頻繁に実行されません。その結果、チャネルc0の占有率およびアクティビティーパーセンテージが高く、チャネルc1の占有率およびアクティビティーの割合が低いことが期待できます。

また、小さなワーク・グループ・サイズを定義すると、占有率が低くなる可能性があり、カーネルが十分なWork-Itemを受け取らない可能性があります。これは問題があります。なぜなら、一般にカーネルの実行中にパイプラインが空であり、パフォーマンスが低下するからです。

メモリーアクセスを確認して、メモリーサイトへのアクセスが連続するメモリーエリアをアドレス指定するように書き換えることができるかどうかを確認します。

通常、ストール率と占有率の合計はほぼ100％になります。ロードおよびストア動作またはチャネルの失速率が高い場合、ロードおよびストア動作またはチャネルがすべてのサイクルを実行できるが、ストールが発生していることを意味します。

グローバルロードおよびストア動作をストールするためのソリューションは以下の通りです。

• ローカルメモリーを使用してデータをキャッシュします。
• データを読み出す回数を減らします。
• グローバル・メモリー・アクセスを改善します。
  • より多くのグローバル・メモリー・フレンドリーなアドレッシング（例えば、ストライド・アクセスからシーケンシャル・アクセスへの変更）のアクセス・パターンを変更してください。
  • カーネルを-no-interleaving=default Intel^® FPGA SDK for OpenCL™オフライン・コンパイラーコマンドオプションでコンパイルし、読み出しバッファーと書き込みバッファーを別々のDDRバンクに分けます。
  • より少ないが広いグローバル・メモリー・アクセスがあります。
• より多くの帯域幅を持つアクセラレーター・ボードを取得します（たとえば、2つのDDRの代わりに3つのDDRを持つボード）。

ローカルロードおよびストア動作をストールするためのソリューションは以下の通りです。

• HTMLエリアレポートを確認してローカルメモリーの構成を確認し、構成を変更してストールフリーにします。

チャネルをストールするためのソリューションは以下の通りです。

• チャンネルの反対側のストールを修正します。たとえば、チャネルの読み出しが停止した場合、チャネルへのライターがチャネルにデータを書き込む速度が十分でなく、調整が必要であることを意味します。
• デザインにチャンネルループがある場合、チャンネルの深さを指定します。

図 71. OpenCLカーネルとプロファイラ解析の例

この例では、dst[]は、 FACTOR2ループの20回の反復とFACTOR1ループの回4回の繰り返しが実行されます。したがって、FACTOR2ループがボトルネックの原因となります。

ループボトルネックを解決するソリューションは以下の通りです。

• FACTOR1をアンロールし、 FACTOR2は均等にループします。 FACTOR1ループをさらに展開するだけでは、ボトルネックは解消されません
• カーネルをベクトル化し、各ループ反復中にMultiple Work-Itemを実行できるようにする

図 72. OpenCLカーネルとプロファイラ解析の例

この例では、アクセラレーター・ボードは25600 MB/sの帯域幅を提供できます。しかし、 vector_addカーネルは、使用可能な帯域幅の14％を要求しています（2リード+ 1ライト）。x 4バイトx 294 MHz = 12バイト/サイクルx 294 MHz = 3528 MB/s。帯域幅を増やすには、各クロックサイクルで実行されるタスクの数を増やします。

低帯域幅のソリューションは以下の通りです。

• カーネルを自動的または手動でベクトル化してwider要求を作成する
• 最も内側のループをアンロールして、クロックサイクルごとにさらに多くの要求を実行する
• タスクの一部を別のカーネルに委譲する

aocl report <filename>.aocx profile.mon <filename>.source

• Profilerは、実行中のカーネルから1組のプロファイル・データのみを抽出することができます。

  Profilerがカーネルの実行が完了した後でプロファイル・データを収集する場合、 profile.monファイルを複数回生成するようにホストAPIを呼び出すことができます。

  カーネルの実行時にプロファイル・データを収集する方法の詳細については、 インテル^® FPGA SDK for OpenCL™ プログラミング・ガイドのカーネル実行部の間に収集プロファイル・データのセクションを参照してください。

• プロファイル・データは、OpenCLプログラムまたは複数のデバイスで永続的ではありません。

  プロファイル・データは、単一のOpenCLプログラムと単一のデバイスのみで要求できます。ホストが新しいカーネルプログラムをFPGAの内部と外部に交換すると、Profilerはプロファイル・データを保存しません。

• パフォーマンス・カウンターを使用してVerilogコードをインスツルメント化すると、ハードウェアリソースの使用率（つまりFPGAエリアの使用率）が増加し、通常はパフォーマンスが低下します。

  パフォーマンス・カウンターを使用したVerilogコードのインストルメントについては、 インテル^® FPGA SDK for OpenCL™ プログラミング・ガイドの「パフォーマンス・カウンターを使用したカーネル・パイプラインのインストルメントのセクションを参照してください。

次のフローチャートは、デザインを反復し、Single Work-Itemカーネルを最適化するために取ることができるアプローチの概要を示しています。使用方法については、 インテル^® FPGA SDK for OpenCL™ EmulatorとProfilerについては、 インテル^® FPGA SDK for OpenCL™ プログラミング・ガイドのOpenCLカーネルのエミュレートとデバッグとOpenCLカーネルの プロファイリングを参照してください。 Intel FPGA Dynamic Profiler for OpenCL GUIおよびプロファイリング情報について詳しくは、 パフォーマンス・ボトルネックを特定するためのカーネルのプロファイルを参照してください。

インテル^® は除去、緩和、単純化、およびローカルメモリーへの転送の順に、単一のWork-Itemカーネルループ搬送依存関係に対処するための以下の最適化オプションを推奨しています。

オフライン・コンパイラーは、コンパイル時および実行時に課される特定の制約に応じて、カーネルのデフォルトのワーク・グループ・サイズを想定しています。

オフライン・コンパイラーは、コンパイル時に次の制約を課します。

• reqd_work_group_size属性に値を指定すると、作業グループのサイズはこの値と一致する必要があります。
• max_work_group_size属性に値を指定すると、ワーク・グループ・サイズはこの値を超えてはなりません。
• reqd_work_group_sizeとmax_work_group_sizeの値を指定せず、カーネルに障壁がある場合、オフライン・コンパイラーのデフォルトの作業グループサイズは256です。
• 両方の属性に値を指定せず、カーネルにバリアが含まれていない場合、オフライン・コンパイラーはコンパイル時にワーク・グループ・サイズに制約を課すことはありません。

OpenCL™標準では、実行時に次の制約が課されます。

• 各ディメンションのワーク・グループ・サイズは、各ディメンションの要求されたNDRangeサイズに均等に分割する必要があります。
• ワーク・グループ・サイズは、 clGetDeviceInfo APIコールのCL_DEVICE_MAX_WORK_GROUP_SIZEおよびCL_DEVICE_MAX_WORK_ITEM_SIZESクエリで指定されたデバイス制約を超えてはなりません。

reqd_work_group_size属性とmax_work_group_size属性の両方に値を指定しない場合、ランタイムはデフォルトのワーク・グループ・サイズを次のように決定します。

• カーネルに障壁が含まれているかローカルのWork-Item IDを参照している場合、またはホストコードでclGetKernelWorkGroupInfoおよびclGetDeviceInfo API呼び出しを使用してワーク・グループ・サイズを照会すると、実行時はワーク・グループ・サイズからWork-Itemにデフォルトで実行します。
• カーネルに障壁がないか、ローカルWork-ItemIDを参照していない場合、またはホストコードがワーク・グループ・サイズを照会しない場合、デフォルトのワーク・グループ・サイズはグローバルNDRangeサイズです。

NDRangeカーネル（つまり、Single Work-Itemカーネルではない）をキューイングするときは、次の条件で明示的なワーク・グループ・サイズを指定します。

• カーネルがメモリーバリア、ローカルメモリー、またはローカルWork-ItemIDを使用している場合。
• ホストプログラムが作業グループのサイズを照会する場合。

カーネルでメモリーバリアが使用されている場合、次のいずれかの作業を実行して、ハードウェア・リソースを最小限に抑えます。

• reqd_work_group_size属性の値を指定します。
• max_work_group_size属性には、すべてのランタイム作業グループサイズ要求に対応する最小の作業グループサイズを割り当てます。

実行時にデフォルトより小さいワーク・グループ・サイズを指定すると、ハードウェアが過剰に消費される可能性があります。したがって、デフォルト以外のワーク・グループ・サイズが必要な場合、 max_work_group_size属性を指定して最大ワーク・グループ・サイズを設定します。すべてのカーネル呼び出しでワーク・グループ・サイズが一定のままである場合、 reqd_work_group_size属性を含めることによって、必要なワーク・グループ・サイズを指定します。reqd_work_group_size属性は、指定したワークグループごとのWork-Itemの数を管理するために、正確なハードウェア量を割り当てるようにオフライン・コンパイラーに指示します。この割り振りにより、ハードウェア・リソースが節約され、カーネルコンピューティング・ユニットの実装効率が向上します。 reqd_work_group_size属性を指定することにより、オフライン・コンパイラーが未知のサイズの作業グループをサポートするために追加のハードウェアを実装するのを防ぐこともできます。

たとえば、次のコードでは、ワークグループの固定サイズを64個のWork-Itemに割り当てることができます。

__attribute__((reqd_work_group_size(64,1,1)))
__kernel void sum (__global const float * restrict a,
                   __global const float * restrict b,
                   __global float * restrict answer)
{
  size_t gid = get_global_id(0);

  answer[gid] = a[gid] + b[gid];
}

カーネルの本体を変更せずに、Work-Itemごとに追加の実行をオフライン・コンパイラーに指示するには、カーネルコードにnum_simd_work_items属性を含めます。次のコードは、ベクトル化係数4を元のカーネルコードに適用します。

__attribute__((num_simd_work_items(4)))
__attribute__((reqd_work_group_size(64,1,1)))
__kernel void sum (__global const float * restrict a,
                   __global const float * restrict b,
                   __global float * restrict answer)
{
   size_t gid = get_global_id(0);

   answer[gid] = a[gid] + b[gid];
}

num_simd_work_items属性を使用するには、reqd_work_group_size属性を使用してカーネルの必要なワーク・グループ・サイズを指定する必要があります。reqd_work_group_sizeに指定したワークグループのサイズは、num_simd_work_itemsに割り当てる値で割り切れなければなりません。上記のコード例では、カーネルの作業グループサイズは64であり、ワークグループは固定されています。各ワークグループ内で、Work-Itemは4つのSIMDベクタレーンに均等に分散されます。オフライン・コンパイラーが4つのSIMDベクタレーンを実装した後、各Work-Itemは4倍の作業を実行するようになりました。

オフライン・コンパイラーはコードをベクトル化し、メモリーアクセスを結合する可能性があります。オフライン・コンパイラーはこれらの最適化を自動的に適用するため、カーネルコードまたはホストコードを変更する必要はありません。

カーネルコードを手動でベクトル化できますが、実装するベクトル化の量を反映するように、ホスト・アプリケーションのNDRangeを調整する必要があります。次の例は、カーネル内で動作を手動で複製するときのコードの変更を示しています。

__kernel void sum (__global const float * restrict a,
                   __global const float * restrict b,
                   __global float * restrict answer)
{
   size_t gid = get_global_id(0);

   answer[gid * 4 + 0] = a[gid * 4 + 0] + b[gid * 4 + 0];
   answer[gid * 4 + 1] = a[gid * 4 + 1] + b[gid * 4 + 1];
   answer[gid * 4 + 2] = a[gid * 4 + 2] + b[gid * 4 + 2];
   answer[gid * 4 + 3] = a[gid * 4 + 3] + b[gid * 4 + 3];
}

この形式では、カーネルは配列aとbから4つの要素を読み出し、合計を計算し、その結果を配列のanswerに格納します 。 FPGAパイプラインはメモリー内の隣接する場所にデータをロードして格納するため、手動でオフライン・コンパイラーに4つのロードおよびストア動作の各グループを結合させることができます。

次の図は、カーネルのパフォーマンスがスタティック・メモリー統合から利益を得られる一般的なケースを示しています。

次のベクトル化されたカーネルを検討してください。

__attribute__((num_simd_work_items(4)))
__attribute__((reqd_work_group_size(64,1,1)))
__kernel void sum (__global const float * restrict a,
                   __global const float * restrict b,
                   __global float * restrict answer)
{
   size_t gid = get_global_id(0);

   answer[gid] = a[gid] + b[gid];
}

OpenCL™カーネルは、メモリー内の連続した場所にアクセスする4つのロード動作を実行します。競合する場所に4回のメモリーアクセスを実行する代わりに、オフライン・コンパイラーは4つの負荷を単一のより広いベクトル負荷に統合します。この最適化により、メモリーシステムへのアクセス回数が減少し、メモリー・アクセス・パターンが改善される可能性があります。

オフライン・コンパイラーは、カーネルをベクトル化する際にスタティック・メモリーの結合を自動的に実行しますが、効率的なメモリーアクセスを確保するために、可能な場合はいつでも、広いベクトルロードとストアをOpenCLコードに使用する必要があります。スタティック・メモリー統合を手動で実装するには、コンパイル時に順次アクセスパターンを識別できるようにコードを記述する必要があります。上の図に示されている元のカーネルコードは、スタティック・メモリー結合から恩恵を受けることができます。これは、バッファーaとbのすべてのインデックスが、コンパイル時に判明しているオフセットで増分するためです。これとは対照的に、次のコードでは、スタティック・メモリーの結合を行うことはできません。

__kernel void test (__global float * restrict a,
		          __global float * restrict b,
                    __global float * restrict answer;
		          __global int * restrict offsets)
{
 size_t gid = get_global_id(0);

 answer[gid*4 + 0] = a[gid*4 + 0 + offsets[gid]] + b[gid*4 + 0];
 answer[gid*4 + 1] = a[gid*4 + 1 + offsets[gid]] + b[gid*4 + 1];
 answer[gid*4 + 2] = a[gid*4 + 2 + offsets[gid]] + b[gid*4 + 2];
 answer[gid*4 + 3] = a[gid*4 + 3 + offsets[gid]] + b[gid*4 + 3];
}

値のオフセット[gid]はコンパイル時には不明です。その結果、オフライン・コンパイラーは静的に読み出しをバッファーにアクセスを合体することはできません。

全体のカーネルスループットを向上させるために、FPGAのハードウェア・スケジューラーはワークグループを追加の使用可能なコンピューティング・ユニットにディスパッチします。コンピューティング・ユニットは、フル・キャパシティーに達していない限り、ワークグループ割り当てに使用できます。

各作業グループは、実行を完了するのに同じ時間がかかるものとします。オフライン・コンパイラーが2つのコンピューティング・ユニットを実装する場合、各コンピューティング・ユニットは作業グループの半分を実行します。ハードウェア・スケジューラーがワークグループをディスパッチするため、このプロセスを独自のコードで管理する必要はありません。

オフライン・コンパイラーは、カーネルの最適な計算単位数を自動的に決定しません。カーネル実装のコンピューティング・ユニットの数を増やすには、以下のコードサンプルに示すように、 num_compute_units属性を使用してオフライン・コンパイラーが作成するコンピューティング・ユニットの数を指定する必要があります。

__attribute__((num_compute_units(2)))
__kernel void sum (__global const float * restrict a,
                   __global const float * restrict b,
                   __global float * restrict answer)
{
    size_t gid = get_global_id(0);

    answer[gid] = a[gid] + b[gid];
}

コンピューティング・ユニットの数を増やすと、スループットが向上します。ただし、以下の図に示すように、コンピューティング・ユニット間のグローバルメモリー帯域幅が増加するという犠牲を払って実行します。また、ハードウェア・リソースの使用率も向上します。

num_compute_units属性とnum_simd_work_items属性の両方は、 Intel^® FPGA SDK for OpenCL™オフライン・コンパイラーがカーネルを実装するために使用するハードウェアの量を増やすことによってスループットを向上させます。num_compute_units属性は、ワークグループのスケジュールを設定できるコンピューティング・ユニットの数を変更し、カーネルがグローバルメモリーにアクセスする回数も変更します。対照的に、 num_simd_work_items属性は、コンピューティング・ユニットが単一のワークグループで並列に実行できる作業量を変更します。 num_simd_work_items属性は、各SIMD ベクタレーンでコントロールロジックを共有することによって、コンピューティング・ユニットのデータパスのみを複製します。

通常、num_simd_work_items属性を使用すると、num_compute_units属性を使用して同じ目標を達成するより効率的なハードウェアにつながります。num_simd_work_items属性を使用すると、オフライン・コンパイラーでメモリーアクセスを結合することもできます。

大域メモリーと競合する複数のコンピューティング・ユニッは、望ましくないメモリー・アクセス・パターンにつながる可能性があります。num_simd_work_itemsが num_compute_units属性の代わりに属性を導入することにより、望ましくないメモリー・アクセス・パターンを変更することができます。また、num_simd_work_itemsは、潜在的に属性num_compute_unitsがオファーを属性同等のカーネル演算ユニットの重複と同じ演算スループットを提供しています。

次のような状況では、 num_simd_work_items属性をカーネルに実装することはできません。

• num_simd_work_itemsに指定する値は、 2,4,8または16ではありません。
• reqd_work_group_sizeの値はnum_simd_work_itemsで割り切れません。

  たとえば、50が4で割り切れないため、次の宣言は正しくありません。

  __attribute__((num_simd_work_items(4)))
  __attribute__((reqd_work_group_size(50,0,0)))

• 複雑な制御フローを持つカーネル。別のWork-Itemは、（例えば、制御パスがget_global_IDまたはget_local_IDに依存）異なる制御経路をたどるするカーネルをベクトル化することはできません。

カーネルコンパイル時に、オフライン・コンパイラーは、ベクトル化最適化の実装が成功したかどうかを通知するメッセージを発行します。報告されたベクトル化係数がnum_simd_work_items属性に指定した値と一致すると、カーネルのベクトル化は成功します。

num_simd_work_items属性が16に設定されたカーネルがFPGAに収まらないケースを検討してください。より狭いSIMDカーネルコンピューティング・ユニットを複製することによってカーネルを修正すると、カーネルが適合するかもしれません。コンピューティング・ユニットの数とSIMD幅の最適なバランスを決定するには、いくつかの実験が必要になるかもしれません。たとえば、4レーン幅のSIMDカーネルコンピューティング・ユニットを3回複製すると、8レーン幅のSIMDカーネルコンピューティング・ユニットを2回複製するよりもスループットが向上する可能性があります。

次のコード例は、 OpenCL™コードでnum_compute_unitsおよびnum_simd_work_items属性を組み合わせる方法を示しています。

__attribute__((num_simd_work_items(4)))
__attribute__((num_compute_units(3)))
__attribute__((reqd_work_group_size(8,8,1)))
__kernel void matrixMult(__global float * restrict C,
                         __global float * restrict A,
. . .

下の図は、上記のカーネルのデータフローを示しています。 num_compute_unitsは、3つの複製コンピューティング・ユニットを実装します。 num_simd_work_itemsは、4つのSIMD ベクタレーンを実装しています。

相互接続トポロジーは、共有されたグローバル、定数、およびローカル・メモリー・システムをそれらの基礎となるメモリーに接続します。相互接続はメモリーポートへのアクセスアービトレーションを含みます。

メモリーアクセスは、共有メモリーリソース（つまり、グローバル、ローカル、および定数メモリー）を競合します。 OpenCLカーネルが多数のメモリーアクセスを実行する場合、 Intel^® FPGA SDK for OpenCL™オフライン・コンパイラーはメモリーアクセス要求を処理するための複雑なアービトレーション・ロジックを生成する必要があります。複雑なアービトレーションロジックにより、最大動作周波数（f _max ）が低下し、カーネルの性能が低下する可能性があります。

以下のセクションでは、メモリーアクセスの最適化について詳しく説明します。要約すると、グローバル・メモリー・アクセスを最小限にすることは、以下の理由から有益である。

• 一般に、OpenCLカーネルのパフォーマンスが向上すると、グローバルメモリー帯域幅要件が増加します。
• グローバルメモリーの最大帯域幅は、最大ローカルメモリー帯域幅よりもずっと小さくなります。
• FPGAの最大計算帯域幅は、グローバルメモリー帯域幅よりもはるかに大きくなります。
  重要: 可能であれば、ローカル、プライベート、または定数メモリーを使用して、カーネルのメモリー帯域幅を増やしてください。

ほとんどの場合、デフォルトのバースト・インターリーブ構成により、メモリーバンク間のロードバランシングが最適化されます。ただし、ロードバランシングを改善するために、2つのインターリーブされていない（および連続した）メモリーエリアとして手動でバンクを分割することが必要な場合もあります。

下の図は、バースト・インターリーブのメモリー・パーティションとインターリーブされていないメモリー・パーティションのメモリー・マッピング・パターンの違いを示しています。

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

__kernel void sum ( __global const float * restrict a,
                    __global const float * restrict b,
                    __global float * restrict c )
{
	size_t gid = get_global_id(0);

	c[gid] = a[gid] + b[gid];
}

アレイaからのロード動作は、Work-ItemのグローバルIDの直接関数であるインデックスを使用します。配列インデックスをWork-ItemのグローバルIDに基づいて設定することにより、オフライン・コンパイラーは連続したロード動作を指示できます。これらのロード動作は、入力配列から順番にデータを検索し、必要に応じてパイプラインに読み取ったデータを送ります。次に、連続ストア動作は、計算パイプラインを出る結果の要素をグローバルメモリー内のシーケンシャルな場所に格納します。

次の図は、連続したメモリーアクセス最適化の例を示しています。

連続したロードおよびストア動作は、アクセス速度の向上およびハードウェア・リソースの必要性の低減につながるため、メモリーアクセス効率を向上させます。データは、パイプラインの計算部分に同時に出入りし、計算とメモリーアクセスの間に重複が可能です。可能な場合、グローバルメモリーにアクセスするロードおよびストア動作の連続するメモリー位置をインデックスするWork-ItemIDを使用します。グローバルメモリーへの順次アクセスは、理想的なアクセスパターンを提供するため、メモリー効率を向上させます。

FPGAボードが異機種グローバル・メモリー・タイプを提供している場合、様々な効率で異なるメモリー・アクセスを処理することに注意してください。

例：

• 長いシーケンシャルアクセスにはDDR SDRAMを使用してください。
• ランダムアクセスにはQDR SDRAMを使用してください。
• ランダムな低レイテンシーアクセスには内蔵RAMを使用してください。

グローバルメモリーにバッファーを割り当てる方法と、異機種バッファーを使用するようにホスト・アプリケーションを変更する方法については、 インテル^® FPGA SDK for OpenCL™ プログラミング・ガイドのグローバルメモリーのバッファーエリアの指定およびグローバルメモリーの手動パーティショニングのためのOpenCLバッファーの割り当てを参照してください。

グローバル・メモリー・アクセスを最小限に抑えるには、まずグローバルメモリーから定数、ローカル、またはプライベート・メモリーへの計算グループからデータをプリロードする必要があります。プリロードされたデータに対してカーネル計算を実行し、その結果をグローバルメモリーに書き戻します。

デフォルトでは、定数キャッシュサイズは16 KBです。定数キャッシュサイズを指定するには、 -const-cache-bytes = <N>aocコマンドのオプション<N>バイト単位の一定のキャッシュサイズです。

長いメモリーレイテンシーを許容するための余分なハードウェアを持つグローバル・メモリー・アクセスとは異なり、定数キャッシュはキャッシュミスに対して大きなパフォーマンス上の不利益を被ります。 OpenCL™カーネルコードの__constant引数がキャッシュに収まらない場合、代わりに__global const引数を使用するとパフォーマンスが向上する可能性があります。ホスト・アプリケーションが定数キャッシュに既にロードされている定数メモリーに書き込む場合、キャッシュされたデータは定数キャッシュから破棄されます（つまり、無効化される）。

-const-cache-bytes=<N> のオプションについて詳しくは インテル^® FPGA SDK for OpenCL™ プログラミング・ガイドの定数メモリーキャッシュサイズの設定 (-const-cache-bytes=<N>)のセクションを参照してください。

レジスターを使用したプライベート・メモリーの実装の詳細については、 インテル^® FPGA SDK for OpenCL™ プログラミング・ガイドのレジスターの推論のセクションを参照してください。

次のコード例は、単一のバンクに実装された8 x 4ローカル・メモリー・システムを示しています。その結果、システム内の2つの要素に並列にアクセスすることはできません。

local int lmem[8][4];

#pragma unroll
for(int i = 0; i<4; i+=2) {
    lmem[i][x] = …; 
}

パフォーマンスを向上させるために、 numbanks（ N ）とbankwidth（ M ）をコードに追加して、メモリーバンクの数とバンク幅をバイト単位で定義することができます。次のコードは、それぞれ16バイト幅の8つのメモリーバンクを実装しています。このメモリーバンク構成により、8×4アレイのパラレルメモリーアクセスが可能になります。

local int __attribute__((numbanks(8), 
                        bankwidth(16)))
                        lmem[8][4]; 
#pragma unroll
for (int i = 0; i < 4; i+=2) {
    lmem[i][x & 0x3] = …; 
}

numbanks（ N ）およびbankwidth（ M ）カーネル属性に異なる値を指定することで、パラレル・アクセス・パターンを変更できます。次のコードは、それぞれ4バイト幅の4つのメモリーバンクを実装しています。このメモリーバンク構成は、8×4アレイ全体にわたる並列メモリーアクセスをイネーブルします。

local int __attribute__((numbanks(4), 
                        bankwidth(4)))
                        lmem[8][4]; 

#pragma unroll
for (int i = 0; i < 4; i+=2) {
    lmem[x][i] = …; 
}

次のコード例は、 numbanksおよびbankwidthに割り当てた値に基づいて、バンクのジオメトリがどのように変化するかを示しています。

local int
__attribute__((numbanks(2),
               bankwidth(16)))
               lmem[2][4]; 

local int 
__attribute__((numbanks(2), 
               bankwidth(8))) 
               lmem[2][4]; 

local int 
__attribute__((numbanks(2), 
               bankwidth(4))) 
               lmem[2][4]; 

local int 
__attribute__((numbanks(4), 
               bankwidth(8))) 
               lmem[2][4]; 

local int
__attribute__((numbanks(4),
               bankwidth(4)))
               lmem[2][4]; 

インテル^® M20Kメモリーブロックには2つのphysicalポートがあります。各M20Kブロックで使用できる論理ポートの数は、ポンピングの程度によって異なります。ポンピングは、他のデザインと比較してM20Kブロックのクロック周波数の尺度です。

カーネルがローカル・メモリー・システムlmemのための3つの読み出しポートと1つの書き込みポートを指定するデザイン例を検討してください。以下のコード例に示すように、ローカル変数宣言のsinglepumpカーネル属性を含めて、M20Kブロックは残りのデザインと同じ頻度で実行されます。

int __attribute__((memory,
                   numbanks(1),
                   bankwidth(64),
                   singlepump,
                   numreadports(3), 
                   numwriteports(1))) 
                   lmem[16];

各シングルポンプM20Kブロックには2つのlogicaポートがあります。ローカル・メモリー・システムの各書き込みポートは、デザインがメモリーシステムを実装するために使用するすべてのM20Kブロックに接続する必要があります。ローカル・メモリー・システムの各読み出しポートは、1つのM20Kブロックに接続する必要があります。これらの接続の制約のため、 lmemに指定されたポート数を実装するには、3つのM20Kブロックが必要です 。

ローカル変数宣言にdoublepumpカーネル属性を含める場合、残りのデザインと同じ頻度でM20Kメモリーブロックを実行するように指定します。

int __attribute__((memory,
                   numbanks(1),
                   bankwidth(64),
                   doublepump,
                   numreadports(3), 
                   numwriteports(1))) 
                   lmem[16];

各ダブルポンプM20Kブロックには4つの論理ポートがあります。そのため、3つの読み出しポートと1つの書き込みポートをすべてlmemに実装するには、1つのM20Kブロックが必要です 。

カーネルのパフォーマンスを最適化するには、通常、追加のFPGAリソースが必要です。これとは対照的に、エリアの最適化はパフォーマンスの低下を招くことがあります。カーネル最適化の間、 インテル^® は、最適なサイズとパフォーマンスのトレードオフを生成するカーネル・プログラミング戦略を決定するために、FPGAボード上で複数のバージョンのカーネルを実行することを推奨します。

1. 画面上の推定資源使用量の要約を確認するには、 aocコマンドに-reportフラグを含めてカーネルをコンパイルします。カーネル固有のエリア使用量の情報を確認するには、 <your_kernel_filename> /reports/report.htmlファイル。
2. 可能であれば、 aocコマンドの-fpcまたは-fp-relaxedオプションを使用してOpenCLカーネルをコンパイルし、浮動小数点計算を実行します。

-report 、 -fp-relaxedおよび-fpcオプションの詳細な使用方法については、 見積もりリソース使用量の要約を表示する（-report）」 、 浮動小数点演算の順序を緩和する（-fp-relaxed）」 、および 浮動小数点数を減らす動作（-fpc）のセクション インテル^® FPGA SDK for OpenCL™ プログラミング・ガイド

浮動小数点演算の詳細については、浮動小数点演算の最適化を参照してください。

たとえば、カーネルに1つの外部メモリーバンクが必要な場合、単一の外部メモリーバンクのみをサポートするボードバリアントをターゲットにします。複数の外部メモリーバンクを持つボードをターゲットにすると、カーネルのエリア使用量が不必要に増加します。

カスタム・プラットフォームでニーズに合ったボード・バリエーションが提供されない場合、ボード・バリエーションの作成を検討してください。詳細については、 カスタム・プラットフォームツールキットユーザーガイドを参照してください。

1. 外部メモリーへのアクセスポイントの数を最小限に抑えます。

   可能であれば、ある場所から入力を読み出し、内部的にデータを処理し、出力を別の場所に書き込むようにカーネルを構造化します。

2. ローカルまたはグローバル・メモリー・アクセスに頼るのではなく、可能であればシフトレジスター推論を使ってカーネルをSingle Work-Itemとして構造化してください。
3. 外部メモリーにデータを書き込むカーネルと外部メモリーからデータを読み込むカーネルを作成する代わりに、直接データ転送のためにカーネル間に インテル^® FPGA SDK for OpenCL™ チャネル拡張を実装します。
4. OpenCLアプリケーションは、多くの独立した定数データアクセスが含まれている場合、代わりに__global constの__constant使用して対応するポインターを宣言します。 __global constを使用する宣言は、ロードまたはストア動作ごとにプライベート・キャッシュを作成します。一方、 __constantを使用した宣言では、単一の定数キャッシュがチップ上にのみ作成されます。
   注意:
   カーネルがCyclone^® Vデバイス（たとえば、Cyclone V SoC）をターゲットにしている場合、 __constantポインターカーネル引数を宣言すると、FPGAのパフォーマンスが低下する可能性があります。

5. カーネルが少数の定数引数を渡した場合、それらをグローバルメモリーへのポインターではなく値として渡します。

   例えば、代わり* COEF __constant INTを通過した後、10にインデックス0を有するCOEFを間接参照の値（INT16のCOEF）としてCOEFを渡します。 coefが__constantポインターの唯一の引数だった場合、それを値として渡すと、定数キャッシュとそれに対応するロードとストア動作が完全に削除されます。

6. パイプライン・ループ内で大規模なシフトレジスターを条件付きでシフトすると、効率の悪いハードウェアが作成されます。たとえば、 if（K> 5）条件が存在する場合 、次のカーネルはより多くのリソースを消費します。

   #define SHIFT_REG_LEN 1024
   __kernel void bad_shift_reg (__global int * restrict src,
                                __global int * restrict dst,
                                int K)
   {
       float shift_reg[SHIFT_REG_LEN];
       int sum = 0;
    
       for (unsigned i = 0; i < K; i++)
       {
           sum += shift_reg[0];
           shift_reg[SHIFT_REG_LEN-1] = src[i];
   
           // This condition will cause sever area bloat.
           if (K > 5)
           {
             #pragma unroll
             for (int m = 0; m < SHIFT_REG_LEN-1 ; m++)
             {
                 shift_reg[m] = shift_reg[m + 1];
             }
           }
           dst[i] = sum;
       }
   }

   重要: 条件付きでシフトレジスターにアクセスしても、ハードウェアの効率は低下しません。 カーネルに大きなシフトレジスターの条件付きシフトを実装する必要がある場合は、ローカルメモリーを使用するようにコードを変更することを検討してください。

1. 必要なときにだけ浮動小数点演算を導入します。
2. Intel^® FPGA SDK for OpenCL™オフライン・コンパイラーは、浮動小数点定数はデータ型を2倍にするようにデフォルト設定されています。 f指定を定数に追加して、単精度浮動小数点演算にします。

   例えば、算術演算sin（1.0）は、倍精度浮動小数点サイン関数を表します。算術演算sin（1.0f）は、単精度浮動小数点正弦関数を表します。

3. 複雑な関数に対して完全精度の結果を必要としない場合、より簡単な算術演算を計算して結果を近似します。以下のシナリオ例を検討してください。
   1. 関数pow（x、n）を計算する代わりに、 nが小さい値の場合、ハードウェア・リソースと面積が大幅に少なくて済むので、繰り返し二乗演算を実行して結果を近似します。
   2. 近似を使用して結果を計算すると、エリアの使用量が過剰になることがあるため、元のエリアと近似されたエリアの使用を認識していることを確認してください。たとえば、 sqrt関数はリソースを消費しません。大まかな近似以外に、実行時にホストが計算しなければならない算術演算でsqrt関数を置き換えると、エリアの使用量が大きくなる可能性があります。
   3. 少数の入力値で作業する場合、代わりにLUTを使用することを検討してください。

4. コンパイル時にオフライン・コンパイラーが計算する定数（たとえばlog（PI/2.0） ）を使用してカーネルが複雑な算術演算を実行する場合、代わりにホスト上で算術演算を実行し、ランタイムでその結果を引数としてカーネルに渡します。

1. アプリケーションに最適なデータ型を選択します。

   たとえば、データ型shortが十分な場合、変数をfloatとして定義しないでください。

2. 算術式の両側が同じデータ型に属していることを確認してください。

   算術式の一方の側が浮動小数点値で、もう一方の側が整数である例を検討してください。一致しないデータ型は、 Intel^® FPGA SDK for OpenCL™オフライン・コンパイラーが暗黙的な変換演算子を作成します。変換演算子は、多数存在する場合は高価になります。

3. パディングがデータ構造内に存在する場合、パディングを使用します。

   たとえば、 float4と同じサイズのfloat3データ型のみが必要な場合、データ型をfloat4に変更して、余分なディメンションを使用して関連のない値を使用することができます。