説明
ワーク アイテムをパイプライン処理してレイテンシおよびスループットを向上します。ワーク アイテムのパイプライン処理は、カーネルのワーク グループに対するループのパイプライン処理の拡張です。これは、カーネルのスループットおよびパフォーマンスを最大にするために必要です。
構文
OpenCL API ソースのパイプライン処理する要素の前に配置します。
__attribute__((xcl_pipeline_workitems))例 1
次の例の reqd_work_group_size 属性を処理するため、Vitis テクノロジにより ND 範囲 (3,1,1) の 3 次元特性を処理するループ ネストが自動的に挿入されます。この追加のループ ネストのため、このカーネルの実行プロファイルはパイプライン処理されていないループのようになります。XCL_PIPELINE_WORKITEMS 属性を追加すると、同時実行性が追加され、コードのスループットが向上します。
kernel
__attribute__ ((reqd_work_group_size(3,1,1)))
void foo(...)
{
...
__attribute__((xcl_pipeline_workitems)) {
int tid = get_global_id(0);
op_Read(tid);
op_Compute(tid);
op_Write(tid);
}
...
}例 2
次の例では、カーネルの適切な要素にワーク アイテム パイプライン処理を追加しています。
__kernel __attribute__ ((reqd_work_group_size(8, 8, 1)))
void madd(__global int* a, __global int* b, __global int* output)
{
int rank = get_local_size(0);
__local unsigned int bufa[64];
__local unsigned int bufb[64];
__attribute__((xcl_pipeline_workitems)) {
int x = get_local_id(0);
int y = get_local_id(1);
bufa[x*rank + y] = a[x*rank + y];
bufb[x*rank + y] = b[x*rank + y];
}
barrier(CLK_LOCAL_MEM_FENCE);
__attribute__((xcl_pipeline_workitems)) {
int index = get_local_id(1)*rank + get_local_id(0);
output[index] = bufa[index] + bufb[index];
}
}