OpenCL カーネルのデータ幅の設定 - 2020.2 Japanese

Vitis 統合ソフトウェア プラットフォームの資料: アプリケーション アクセラレーション開発 (UG1393)

Document ID
UG1393
Release Date
2021-03-22
Version
2020.2 Japanese

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)))
ワーク アイテムの合計を定義します (XYZ は正の定数)。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 は展開されたループで計算される (反復ごとに乗算されたスループットが使用される) ので、パフォーマンスを落としてまで直す必要はありません (通常はそれでもパフォーマンスは向上します)。
    • リソース制約がある可能性があり、これらの問題を解決するとパフォーマンスがさらに改善することを示すメッセージがスケジューラから表示されます。
    • ローカル配列は、通常ベクター化しない方法をでコードの後のセクションでアクセスされるので、自動的に再形成されません。