2.8.7. ロード・ストア・ユニット

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

ダウンロード PDF

ID 683521

日付 12/08/2017

バージョン 17.1

Public

2.8.7. ロード・ストア・ユニット

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

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

ロード・ストア・ユニットのタイプ

コンパイラーは、推論されたメモリー・アクセス・パターン、ターゲットプラットフォームで使用可能なメモリーの種類、およびメモリーアクセスがローカルメモリーかグローバルメモリーかに基づいて、いくつかの異なる種類のロード・ストア・ユニット（LSU）を生成できます。 Intel® FPGA SDK for OpenCL™オフライン・コンパイラーは、次のタイプのLSUを生成できます。

バースト合体のロード・ストア・ユニット

バースト合体LSUは、コンパイラーによってインスタンス化されるデフォルトのLSUタイプです。可能な限り大きなバーストが生成されるまで要求をバッファーします。バースト合体LSUは大域メモリーへの効率的なアクセスを提供できますが、相当量のFPGAリソースが必要です。


kernel void burst_coalesced (global int * restrict in, 
                             global int * restrict out) {
  int i = get_global_id(0);
  int value = in[i/2];      // Burst-coalesced LSU
  out[i] = value;
}

メモリー・アクセス・パターンやその他の属性によっては、次の方法でバースト合体LSUを変更することがあります。

ロード・ストア・ユニットのプリフェッチ

先読みLSUは、先行するアドレスに基づいてFIFOに有効なデータを完全に保持し、連続した読み出しを仮定するために、バーストがメモリーから大きなブロックを読み出すFIFO（時には名前付きパイプと呼ばれる）をインスタンス化します。不連続リードはサポートされていますが、FIFOをフラッシュして再充填する際に不利益が生じます。


kernel void prefetching (global int * restrict in, 
                         global int * restrict out, 
                         int N) {
  int res = 1;
  for (int i = 0; i < N; i++) {
    int v = in[i];             // Prefetching LSU
    res ^= v;
  }
  out[0] = res;
}

ストリーミングロードストアユニット

ストリーミングLSUは、FIFOが有効なデータでいっぱいになるように、大きなブロックをメモリーからFFIFOをインスタンス化します。このデータブロックは、メモリーアクセスが順序通りであり、アドレスがベースアドレスからの単純なオフセットとして計算できる場合にのみ使用できます。


kernel void streaming (global int * restrict in, 
                       global int * restrict out) {
  int i = get_global_id(0);
  int idx = out[i];           // Streaming LSU
  int cached_value = in[idx];
  out[i] = cached_value;      // Streaming LSU
}

セミストリーミングロードストアユニット

セミストリーミングLSUは、読み出し専用キャッシュをインスタンス化します。キャッシュにはエリアのオーバーヘッドがありますが、グローバルメモリー内の同じデータ位置に繰り返しアクセスする場合、パフォーマンスが向上します。カーネル内のストアによってデータが上書きされないようにする必要があります。これは、キャッシュの一貫性を損なうためです。 LSUキャッシュは、関連するカーネルが開始されるたびにフラッシュされます。


#define N 16
kernel void semi_streaming (global int * restrict in, 
                            global int * restrict out) {
  #pragma unroll 1
  for (int i = 0; i < N; i++) {
    int value = in[i]; // Semi-streaming LSU
    out[i] = value;
  }
}

ローカル・パイプラインロードストアユニット

ローカル・パイプライン化されたLSUは、ローカルメモリーにアクセスするために使用されるパイプライン化されたLSUです。リクエストは受信するとすぐに提出されます。メモリーアクセスはパイプライン化されているので、一度に複数のリクエストを飛行することができます。 LSUとローカルメモリーとの間にアービトレーションがない場合、ローカル・パイプライン化されたノーストールLSUが作成されます。


__attribute((reqd_work_group_size(1024,1,1)))
kernel void local_pipelined (global int* restrict in, 
                             global int* restrict out) {
  local int lmem[1024];
  int gi = get_global_id(0);
  int li = get_local_id(0);

  int res = in[gi];
  for (int i = 0; i < 4; i++) {
    lmem[li - i] = res;                 // Local-pipelined LSU
    res >>= 1;
  }

  barrier(CLK_GLOBAL_MEM_FENCE);

  res = 0;
  for (int i = 0; i < 4; i++) {
    res ^= lmem[li - i];                // Local-pipelined LSU
  }

  out[gi] = res;
}

コンパイラーは、ローカル・パイプライン化されたLSUを次のように変更する可能性があります。

Never-stall

Global Infrequent Load-Store Units

グローバルな頻度の低いLSUは、まれであることが証明できるグローバル・メモリー・アクセスに使用されるパイプライン型のLSUです。グローバルなまれなLSUは、ループに含まれていないメモリー動作に対してのみインスタンス化され、NDRangeカーネル内の単一のスレッドに対してのみアクティブです。

パイプライン化されたLSUは他のLSUタイプよりも小さいため、コンパイラーはパイプライン化されたLSUとしてグローバルなまれなLSUを実装します。パイプライン化されたLSUのスループットは低下する可能性がありますが、メモリーアクセスがまれであるため、このスループットのトレードオフは許容されます。


kernel void global_infrequent (global int * restrict in, 
                               global int * restrict out, 
                               int N) {
  int a = 0;
  if (get_global_id(0) == 0)
    a = in[0];                   // Global Infrequent LSU
  for (int i = 0; i < N; i++) {
    out[i] = in[i] + a;
  }
}

コンスタント・パイプライン・ロード・ストア・ユニット

一定のパイプライン化されたLSUは、主に定数キャッシュからの読み出しに使用されるパイプライン化されたLSUです。一定のパイプライン化されたLSUは、バースト合体LSUより少ない面積を消費する。一定パイプライン化されたLSUのスループットは、リードが定数キャッシュ内でヒットしたかどうかによって大きく異なります。キャッシュミスは高価です。


 kernel void constant_pipelined (constant int *src, 
                                 global int *dst) {
  int i = get_global_id(0);
  dst[i] = src[i];              // Constant pipelined LSU
}

インスタンスIDについて詳しくは、キャッシュ・メモリーを参照してください。

原子パイプライン式ロード・ストア・ユニット

アトミックパイプライン化されたLSUは、すべてのアトミック動作に使用されます。アトミック動作を使用すると、カーネルのパフォーマンスが大幅に低下する可能


kernel void atomic_pipelined (global int* restrict out) {
  atomic_add(&out[0], 1);  // Atomic LSU
}

ロードストアユニット修飾子

カーネルのメモリー・アクセス・パターンに応じて、コンパイラーはいくつかのLSUを変更します。

キャッシュ

バースト合体LSUにはキャッシュが含まれることがあります。キャッシュは、メモリー・アクセス・パターンがデータ依存であるか、または繰り返しているように見える場合に作成されます。ロードで同じデータが必要な場合でも、キャッシュを他のロードと共有することはできません。キャッシュはカーネル開始時にフラッシュされ、キャッシュなしで同等のLSUより多くのハードウェア・リソースを消費します。キャッシュは、アクセスパターンを簡素化するか、ポインターを揮発性としてマークすることによって無効にすることができます。


kernel void cached (global int * restrict in, 
                    global int * restrict out) {
  int i = get_global_id(0);
  int idx = out[i]; 
  int cached_value = in[idx];  // Burst-coalesced cached LSU
  out[i] = cached_value;
}

ライト・アクノリッジ（ライト・アクノリッジ）

バースト集約されたストアLSUは、データの依存関係が存在する場合、書き込み確認信号を必要とすることがあります。ライトアクノリッジ信号を有するLSUは、追加のハードウェア資源を必要とする。複数のライトアクノリッジLSUが同じメモリーにアクセスすると、スループットが低下する可能性があります。


kernel void write_ack (global int * restrict in, 
                       global int * restrict out, 
                       int N) {
  for (int i = 0; i < N; i++) {
    if (i < 2)
      out[i] = 0;            // Burst-coalesced write-ack LSU
    out[i] = in[i];
  }
}

非整列

バースト合体LSUが外部メモリー・ワード・サイズにアラインメントされていないメモリーにアクセスできる場合、アラインメントされていないLSUが作成されます。アラインメントされていないLSUを実装するには、追加のハードウェア・リソースが必要です。アラインメントされていない多くの要求を受信すると、アラインメントされていないLSUのスループットが低下する可能性があります。


kernel void non_aligned (global int * restrict in, 
                         global int * restrict out) {
  int i = get_global_id(0);
  
  // three loads are statically coalesced into one, creating a Burst-coalesced non-aligned LSU
  int a1 = in[3*i+0];
  int a2 = in[3*i+1];
  int a3 = in[3*i+2];
  
  // three stores statically coalesced into one
  out[3*i+0] = a3;
  out[3*i+1] = a2;
  out[3*i+2] = a1;
}

Never-stall

ローカル・パイプライン化されたLSUがアービトレーションせずにローカルメモリーに接続されている場合、メモリーへのすべてのアクセスがコンパイラーに知られている一定数のサイクルになるため、ストールしないLSUが作成されます。

次の例では、96ビット幅のメモリーアクセスの一部は2つのメモリーワードにまたがるため、メモリーから2つのフルラインのデータを読み出す必要があります。


__attribute((reqd_work_group_size(1024,1,1)))
kernel void never_stall (global int* restrict in, 
                         global int* restrict out, 
                         int N) {
  local int lmem[1024];
  int gi = get_global_id(0);
  int li = get_local_id(0);

  lmem[li] = in[gi];                   // Local-pipelined never-stall LSU
  barrier(CLK_GLOBAL_MEM_FENCE);
  out[gi] = lmem[li] ^ lmem[li + 1];
}

言語の選択

Intel.com サーチを使用

クイックリンク

最近の検索

高度検索

検索のみ

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

2.8.7. ロード・ストア・ユニット

ロード・ストア・ユニットのタイプ

バースト合体のロード・ストア・ユニット

ロード・ストア・ユニットのプリフェッチ

ストリーミングロードストアユニット

セミストリーミングロードストアユニット

ローカル・パイプラインロードストアユニット

Global Infrequent Load-Store Units

コンスタント・パイプライン・ロード・ストア・ユニット

原子パイプライン式ロード・ストア・ユニット

ロードストアユニット修飾子

キャッシュ

ライト・アクノリッジ（ライト・アクノリッジ）

非整列

Never-stall