説明
ループをパイプライン処理して、レイテンシを向上し、カーネル スループットおよびパフォーマンスを最大限にします。
ループを展開することにより同時実行性は増加しますが、カーネル データパスのすべての要素を常にビジー状態に保持する問題は解決されません。ループ展開されていても、ループ制御依存により順次実行となることがあります。演算が順次実行されると、ハードウェアがアイドル状態となり、パフォーマンスが低下します。
ザイリンクスでは、この問題に対処するため、XCL_PIPELINE_LOOP 属性を使用したループ パイプライン処理用に OpenCL 2.0 API 仕様上にベンダー拡張を導入しました。
デフォルトでは、v++ コンパイラにより自動的に、トリップカウントが 64 を超える場合はループがパイプライン処理され、トリップカウントが 64 以下の場合はループが展開されます。これにより、良い結果が得られるはずです。ループの前に NOUNROLL 属性と XCL_PIPELINE_LOOP 属性を指定すると、ループをパイプライン処理するよう指定できます。
構文
OpenCL ソースのループ宣言の前に配置します。
__attribute__((xcl_pipeline_loop(<II_number>)))説明:
- <II_number>: パイプラインの開始間隔 (II) を指定します。Vitis HLS ツールでこの指定を満たすよう試みられますが、データ依存性によって実際の開始間隔 (II) がこれより大きくなる場合があります。II を指定しない場合、デフォルトは 1 です。
例
次の例では、指定した関数に含まれる for ループの II を 3 に指定しています。
__kernel void f(__global int *a) {
__attribute__((xcl_pipeline_loop(3)))
for (unsigned i = 0; i < 64; ++i)
a[i] = i;
}