説明
4K サイズよりも大型のカーネルを指定する必要がある場合は、REQD_WORK_GROUP_SIZE ではなくこの属性を使用してください。
この属性は、Vitis コア開発キットで reqd_work_group_size 属性によりサポートされるデフォルトの最大ワーク グループ サイズを拡大します。XCL_MAX_WORK_GROUP_SIZE 属性を使用すると、Vitis コア開発キットで 4096 より大型のワーク サイズをサポートできます。
注記: 実際のワーク グループ サイズの制限は、プラットフォームに選択したザイリンクス デバイスによって異なります。
構文
カーネル定義の前、またはそのカーネル用に指定されたプライマリ関数の前に配置します。
__attribute__((xcl_max_work_group_size(<X>, <Y>, <Z>)))説明:
- <X>、<Y>、<Z>: カーネルの ND 範囲を指定します。カーネルのワーク グループのサイズを指定する 3 次元行列の各次元を表します。
例
次は、最適化されていない加算器のカーネル ソース コード例です。このデザインでは、ワーク サイズを行列のサイズ (64x64 など) に設定している以外は、属性は指定されていません。つまり、ワーク グループ全体を実行すると、入力行列 a および b が完全に加算され、結果が出力されます。3 つはすべてグローバル整数ポインターであり、行列の各値は 4 バイトで、オフチップの DDR グローバル メモリに格納されます。
#define RANK 64
__kernel __attribute__ ((reqd_work_group_size(RANK, RANK, 1)))
void madd(__global int* a, __global int* b, __global int* output) {
int index = get_local_id(1)*get_local_size(0) + get_local_id(0);
output[index] = a[index] + b[index];
}このローカル ワーク サイズ (64, 64, 1) は、グローバル ワーク サイズと同じです。この設定により、合計ワーク サイズ 4096 が作成されます。
注記: これは、標準の OpenCL 属性 REQD_WORK_GROUP_SIZE を使用した場合に Vitis コア開発キットでサポートされる最大ワーク サイズです。ザイリンクス 属性の
xcl_max_work_group_size を使用すると、Vitis コア開発キットで 4096 より大型のワーク サイズをサポートできます。64x64 より大型の行列では、ワーク サイズを定義するのに 1 次元のみを使用する必要があります。つまり、128x128 行列はワーク サイズ (128, 1, 1) のカーネルで演算できます (各実行でデータの行全体または列全体を演算)。