API には、OpenCL カーネル用に AXI データ幅の増加する属性があります。コードを手動で変更しなくてもよいようにするには、次の OpenCL 属性を使用して、データ幅の拡張とアルゴリズムのベクター化を実行します。
次のような例があるとします。
__attribute__((reqd_work_group_size(64, 1, 1)))
__attribute__((vec_type_hint(int)))
__attribute__((xcl_zero_global_work_offset))
__kernel void vector_add(__global int* c, __global const int* a, __global const int* b) {
size_t idx = get_global_id(0);
c[idx] = a[idx] + b[idx];
}
この場合、ハード コード化されたインターフェイスは 32 ビット幅のデータパス (int *c, int* a, int *b)
なので、直接インプリメントするとメモリ スループットが大幅に制限されますが、3 つの属性の値に基づいて、幅の自動拡張と変換が適用されます。
-
__attribute__((vec_type_hint(int)))
-
int
が計算およびメモリ転送 (32 ビット) に主に使用されるデータ型であることを宣言します。これにより、AXI インターフェイスのターゲット帯域幅 (512 ビット) に基づいて、ベクター化/幅拡張の係数を計算できます。この例では係数は 16 = 512 ビット / 32 ビットで、理論上はベクター化が適用されると 16 個の値が処理されます。 -
__attribute__((reqd_work_group_size(X, Y, Z)))
- ワーク アイテムの合計を定義します (
X
、Y
、Z
は正の定数)。X*Y*Z
はワーク アイテムの最大数であり、メモリ帯域幅が飽和する最大ベクター化係数を定義します。この例の場合、ワーク アイテムの合計は64*1*1=64
です。適用される実際のベクター化係数は、実際にコード記述されたデータ型または
vec_type_hint
で定義されるベクター化係数と、reqd_work_group_size
で定義される可能な最大ベクター係数の最大公約数です。可能な最大ベクター化係数を実際のベクター化係数で除算した商は、OpenCL 記述の残りのループ カウントとなります。このループはパイプライン処理されるので、複数のループ反復が残っている場合に、パイプライン処理されたインプリメンテーションの利点を活かすことができます。特にベクター化された OpenCL コードに長いレイテンシがある場合に有益です。
-
__attribute__((xcl_zero_global_work_offset))
-
__attribute__((xcl_zero_global_work_offset))
は、ランタイムでグローバル オフセット パラメーターは使用されず、すべてのアクセスが揃っていることをコンパイラに指定します。これにより、ワーク グループのアライメントに関する有益な情報がコンパイラに伝わり、通常メモリ アクセスのアライメントに伝搬されます (ハードウェアがより少ない)。
これらの変換により、合成される実際のデザインが変わることに注意してください。部分的に展開されるループの場合、データが格納されるローカル配列の形状を変更する必要があります。これは通常問題なく動作しますが、まれに悪影響のあることがあります。
次に例を示します。
- 配列の分割で、分割係数を展開/ベクター化係数で除算できない。
- このため、マルチプレクサーが多数必要となり、スケジューラで問題となります (メモリ使用量およびコンパイル時間が大幅に増加する可能性あり)。ザイリンクスでは、分割係数を 2 のべき乗にすることをお勧めします (ベクター係数は常に 2 のべき乗)。
- ベクター化されるループに関係のないリソース制約があると、スケジューラで II が満たされないことを示すメッセージが表示される。
- II は展開されたループで計算される (反復ごとに乗算されたスループットが使用される) ので、パフォーマンスを落としてまで直す必要はありません (通常はそれでもパフォーマンスは向上します)。
- リソース制約がある可能性があり、これらの問題を解決するとパフォーマンスがさらに改善することを示すメッセージがスケジューラから表示されます。
- ローカル配列は、通常ベクター化しない方法をでコードの後のセクションでアクセスされるので、自動的に再形成されません。