インテルのみ表示可能 — GUID: mwh1391807503031
Ixiasoft
6.2.1. スタティック・メモリー統合
静的メモリー統合は、カーネルが非プライベート・メモリーにアクセスする回数を減らそうとする Intel® FPGA SDK for OpenCL™オフライン・コンパイラー最適化ステップです。
次の図は、カーネルのパフォーマンスがスタティック・メモリー統合から利益を得られる一般的なケースを示しています。
図 74. スタティック・メモリー統合
次のベクトル化されたカーネルを検討してください。
__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]はコンパイル時には不明です。その結果、オフライン・コンパイラーは静的に読み出しをバッファーにアクセスを合体することはできません。