OpenCL 属性

このセクションでは、Vitis コア開発キットおよび Vivado® HLS ツール合成でのシステム最適化のために、ソース コードに追加可能な OpenCL™ 属性について説明します。

Vitis コア開発キットでは、コードをデータの動きとカーネル パフォーマンスで最適化する OpenCL 属性が提供されています。データの動きの最適化は、インターフェイス帯域幅および DDR 帯域幅の最大限に活用することにより、システム レベルのデータ スループットを最大にすることを目的として実行されます。カーネル計算最適化は、カーネル インターフェイスにデータが到達したらすぐにすべてのデータを消費できるプロセッシング ロジックを作成することを目的として実行されます。これは通常、関数のインライン展開とパイプライン処理、ループ展開、配列分割、データフローなどの手法を使用してデータパスを一致させるようにプロセッシング コードを展開することによって達成されます。

OpenCL 属性には、次の表で指定されるようなタイプがあります。

表 1. OpenCL 属性 (タイプ別)
タイプ 属性
カーネル最適化
関数のインライン展開
タスク レベルのパイプライン処理
パイプライン処理
ループ最適化
配列最適化
注記: 配列変数では、使用できる配列最適化属性は 1 つのみです。
ヒント: Vitis コンパイラでは、gcc でサポートされる次のような多くの標準属性もサポートされます。
  • ALWAYS_INLINE
  • NOINLINE
  • UNROLL
  • NOUNROLL

always_inline

説明

ALWAYS_INLINE 属性は、関数をインライン展開する必要があることを示します。この属性は GCC の標準機能であり、Vitis コンパイラの標準機能でもあります。

ヒント: NOINLINE 属性は GCC の標準機能であり、Vitis コンパイラでもサポートされます。

この属性は、関数を呼び出し関数内でインライン展開するコンパイラ最適化をイネーブルにします。インライン展開された関数は、RTL で別の階層としては表示されなくなります。

関数をインライン展開すると、関数内の演算が共有され、呼び出し関数の周辺の演算と効率よく最適化されることがあります。ただし、インライン展開された関数はほかの関数と共有できなくなるので、インライン展開された関数とその関数の別のインスタンス (より広く共有可能) でロジックが複製される可能性があります。これによりパフォーマンスを向上できますが、RTL をインプリメントするのに必要なエリアが増加します。

OpenCL カーネルの場合、Vitis コンパイラでは独自の規則を使用して、関数がインライン展開または非インライン展開されます。インライン関数を直接制御するには、ALWAYS_INLINE または NOINLINE 属性を使用ます。

デフォルトでは、インライン展開は関数階層のすぐ下の階層でのみ実行され、サブ関数では実行されません。

重要: XCL_DATAFLOW 属性を使用すると、コンパイラで ALWAYS_INLINE 属性が無視され、関数がインライン展開されません。

構文

OpenCL API ソースで関数の前に記述し、その関数が呼び出されたときに常にインライン展開されるようにします。
__attribute__((always_inline))

次の例では、関数 foo に ALWAYS_INLINE 属性を追加しています。

__attribute__((always_inline))
  void foo ( a, b, c, d ) {
  ...
}

次の例では、関数 foo がインライン展開されないようにしています。

__attribute__((noinline))
  void foo ( a, b, c, d ) {
  ...
}

関連項目

opencl_unroll_hint

説明

重要: これはコンパイラ ヒントであり、コンパイラにより無視される可能性があります。

ループ展開は、Vitis コンパイラで使用可能な最適化手法です。ループ展開最適化は、コンパイラで同時処理が認識されるようにするために実行されます。認識された新たな同時処理により、レイテンシが削減され、パフォーマンスが向上しますが、より多くの FPGA ファブリック リソースを使用します。

OPENCL_UNROLL_HINT 属性は OpenCL 仕様の一部であり、Vitis コンパイラでループ (forwhiledo) が展開されるよう指定します。詳細は、ループ展開 を参照してください。

OPENCL_UNROLL_HINT 属性修飾子は、適用するループの直前に記述する必要があります。この属性では、ループの完全な展開、指定した量の部分展開、またはループ展開のディスエーブルを指定できます。

構文

OpenCL ソースのループ宣言の前に配置します。

__attribute__((opencl_unroll_hint(<n>)))

説明:

  • <n>: オプションのループ展開係数。正の整数またはコンパイル時定数表現で指定します。1 にすると、ループ展開がディスエーブルになります。
    ヒント: <n> を指定しない場合、ループの展開係数はコンパイラにより自動的に決定されます。

例 1

次の例では、for ループを係数 2 で展開しています。この結果、計算ユニットの 4 つの順次反復ではなく、2 つの並列ループ反復が生成されます。

__attribute__((opencl_unroll_hint(2)))
for(int i = 0; i < LENGTH; i++) {
bufc[i] = bufa[i] * bufb[i];
}

上記のループは、コンパイラにより次のコードに変換されます。

for(int i = 0; i < LENGTH; i+=2) {
bufc[i] = bufa[i] * bufb[i];
bufc[i+1] = bufa[i+1] * bufb[i+1];
}

関連項目

reqd_work_group_size

説明

OpenCL API カーネルが OpenCL デバイスでの実行用に投稿されると、ND 範囲 (1、2、または 3 次元) と呼ばれるインデックス空間内で実行されます。これは、OpenCL API ではグローバル サイズと呼ばれます。ワーク グループ サイズは、カーネル計算ユニット (CU) の 1 回の起動で処理可能な ND 範囲の量を定義します。ワーク グループ サイズは、OpenCL API ではローカル サイズとも呼ばれます。OpenCL コンパイラでは、カーネルおよび選択されたデバイスのプロパティに基づいてワーク グループ サイズを決定できます。ワーク グループ サイズ (ローカル サイズ) が決定されたら、ND 範囲 (グローバル サイズ) が自動的にワーク グループに分割され、デバイス上で実行するためにワーク グループがスケジューリングされます。

OpenCL コンパイラでワーク グループ サイズを定義できますが、カーネルの FPGA インプリメンテーションでは、カーネルの REQD_WORK_GROUP_SIZE 属性でワーク グループ サイズを定義することをお勧めします。この属性は、カーネルのカスタム ロジックの生成中にパフォーマンスを最適化するのに推奨されます。

ヒント: FPGA インプリメンテーションの場合は、REQD_WORK_GROUP_SIZE 属性をカーネルのカスタム ロジックの生成中にパフォーマンスを最適化するために使用できるので、この属性を指定することをお勧めします。

OpenCL カーネル関数は、ND 範囲インデックス空間の各点に対して 1 回のみ実行されます。ND 範囲の各点に対するこの処理ユニットは、ワーク アイテムと呼ばれます。ワーク アイテムは、計算ユニットにスケジューリングされるワーク ユニットであるワーク グループにまとめられます。オプションの REQD_WORK_GROUP_SIZE 属性は、clEnqueueNDRangeKernellocal_work_size 引数として使用される必要のある計算ユニットのワーク グループ サイズを定義します。これにより、生成されたコードがこのカーネル用に適切に最適化されます。

構文

カーネル定義の前、またはそのカーネル用に指定されたプライマリ関数の前に配置します。

__attribute__((reqd_work_group_size(<X>, <Y>, <Z>)))

説明:

  • <X>、<Y>、<Z>: カーネルの ND 範囲を指定します。カーネルのワーク グループのサイズを指定する 3 次元行列の各次元を表します。

次の OpenCL C カーネル コードはベクトル加法デザインを記述したもので、2 つの配列のデータの和が 3 つ目の配列に挿入されます。ワーク グループに必要なサイズは 16x1x1 です。このカーネルは 16 回実行され、有効な結果が生成されます。

#include <clc.h>
// For VHLS OpenCL C kernels, the full work group is synthesized
__attribute__ ((reqd_work_group_size(16, 1, 1)))
__kernel void 
vadd(__global int* a,
__global int* b,
__global int* c)
{
int idx = get_global_id(0);
c[idx] = a[idx] + b[idx];
}

関連項目

vec_type_hint

説明

重要: これはコンパイラ ヒントであり、コンパイラにより無視される可能性があります。

オプションの __attribute__((vec_type_hint(<type>)))OpenCL 言語仕様の一部であり、カーネルの計算幅を表す OpenCL コンパイラへのヒントです。この設定は、コンパイラでコードを自動ベクター化しようとするときにプロセッサ帯域幅の使用率を計算する際のベースとなります。

デフォルトでは、カーネルに __attribute__((vec_type_hint(int))) 修飾子があるとして処理されます。これにより、異なるベクター化タイプを指定できます。

自動ベクター化では、コンパイラでワーク アイテムが結合または分離される状況に対応するため、カーネルから呼び出されたライブラリはランタイム時に再コンパイル可能であると想定されます。つまり、ライブラリはハード コードされたバイナリではないか、ハード コードされたバイナリの場合はソースまたはリターゲット可能な中間表記も提供されていることが必要です。これが、コード セキュリティの問題となる可能性もあります。

構文

カーネル定義の前、またはそのカーネル用に指定されたプライマリ関数の前に配置します。
__attribute__((vec_type_hint(<type>)))

説明:

  • <type>: 次の表にリストされたビルトイン ベクターのいずれか、または構成スカラー要素のデータ型。
    注記: 指定しない場合は、INT 型であると想定されます。
表 2. ベクター型

タイプ

説明
char<n> <n> 個の 8 ビット符号付き 2 の補数整数値のベクター。
uchar<n> <n> 個の 8 ビット符号なし整数値のベクター。
short<n> <n> 個の 16 ビット符号付き 2 の補数整数値のベクター。
ushort<n> <n> 個の 16 ビット符号なし整数値のベクター。
int<n> <n> 個の 32 ビット符号付き 2 の補数整数値のベクター。
uint<n> <n> 個の 32 ビット符号なし整数値のベクター。
long<n> <n> 個の 64 ビット符号付き 2 の補数整数値のベクター。
ulong<n> <n> 個の 64 ビット符号なし整数値のベクター。
float<n> <n> 個の 32 ビット浮動小数点値のベクター。
double<n> <n> 個の 64 ビット浮動小数点値のベクター。
注記: <n> を指定しない場合は、1 と想定されます。上記のベクター データ型名で <n> が 2、3、4、8、および 16 以外の値のものも、予約されています。つまり、<n> は 2、3、4、8、および 16 のみに指定できます。

次の例では、基本計算幅が 2 倍幅整数であると想定して自動ベクター化を実行しています。

#include <clc.h>
// For VHLS OpenCL C kernels, the full work group is synthesized
__attribute__((vec_type_hint(double)))
__attribute__ ((reqd_work_group_size(16, 1, 1)))
__kernel void 
...

関連項目

work_group_size_hint

説明

重要: これはコンパイラ ヒントであり、コンパイラにより無視される可能性があります。

OpenCL API 規格のワーク グループ サイズは、カーネル計算ユニットの 1 回の起動で処理可能な ND 範囲のサイズを定義します。OpenCL カーネルが OpenCL デバイスでの実行用に投稿されると、ND 範囲 (1、2、または 3 次元) と呼ばれるインデックス空間内で実行されます。

OpenCL カーネル関数は、ND 範囲インデックス空間の各点に対して 1 回のみ実行されます。ND 範囲の各点に対するこの処理ユニットは、ワーク アイテムと呼ばれます。ループ反復が順番どおりに順次実行される C の for ループとは異なり、OpenCL ランタイムおよびデバイスではワーク アイテムを並列に任意の順序で実行できます。

ワーク アイテムは、計算ユニットにスケジューリングされるワーク ユニットであるワーク グループにまとめられます。オプションの WORK_GROUP_SIZE_HINT 属性は OpenCL 言語仕様の一部であり、ほとんどの場合に clEnqueueNDRangeKernel への local_work_size 引数により指定されるワーク グループ サイズ値を示す OpenCL コンパイラへのヒントです。これにより、生成されたコードが指定した値に従って最適化されます。

ヒント: FPGA インプリメンテーションの場合、REQD_WORK_GROUP_SIZE の、代わりに WORK_GROUP_SIZE_HINT 属性をカーネルのカスタム ロジックの生成中にパフォーマンスを最適化するために使用できるので、この属性を指定することをお勧めします。

構文

カーネル定義の前、またはそのカーネル用に指定されたプライマリ関数の前に配置します。

__attribute__((work_group_size_hint(<X>, <Y>, <Z>)))

説明:

  • <X>、<Y>、<Z>: カーネルの ND 範囲を指定します。カーネルのワーク グループのサイズを指定する 3 次元行列の各次元を表します。

次の例では、カーネルがほとんどの場合にワーク グループ サイズ 1 で実行されることをコンパイラにヒントとして指定しています。

__attribute__((work_group_size_hint(1, 1, 1)))
__kernel void
...

関連項目

xcl_array_partition

説明

重要: 配列変数では、使用できる属性は 1 つのみです。XCL_ARRAY_PARTITION では多次元配列がサポートされますが、1 つの属性で再形成できるのは 1 つの次元のみです。

OpenCL プログラマのほかの計算デバイスと比較した場合の FPGA の利点は、システム全体のメモリ アーキテクチャをカスタマイズし、計算ユニットに挿入できることです。デフォルトでは、Vitis コンパイラにより、カーネル コードのスタティック コード解析に基づいてローカルおよびプライベート メモリの帯域幅を最大限にする計算ユニット内にメモリ アーキテクチャが生成されます。これらのメモリはカーネル ソース コードの属性に基づいてさらに最適化することが可能であり、ローカルおよびプライベート メモリの物理的なレイアウトおよびインプリメンテーションを指定するのに使用できます。計算ユニットのメモリの物理的なレイアウトを制御する Vitis コンパイラの属性は array_partition です。

1 次元配列の場合、XCL_ARRAY_PARTITION 属性により、カーネル コード内で宣言された配列が 1 つの物理メモリではなく複数の物理メモリにインプリメントされます。どの分割方法を使用するかは、アプリケーションおよびパフォーマンス要件によって異なります。Vitis コンパイラで使用可能な配列の分割方法は、cyclicblock、および complete です。

構文

この属性は、配列変数の定義と一緒に指定します。

__attribute__((xcl_array_partition(<type>, <factor>, 
<dimension>)))

説明:

  • <type>: 次のいずれかの分割タイプを指定します。
    • cyclic: サイクリック分割では、配列が計算ユニットのロジックにより同時にアクセス可能な複数の小型の物理メモリとしてインプリメントされます。各メモリに要素が 1 つずつ配置され、すべての配列に配置されたら最初の配列に戻って、配列が完全に分割されるまでそれが繰り返されます。
    • block: ブロック分割では、配列が計算ユニットのロジックにより同時にアクセス可能な複数の小型メモリとしてインプリメントされます。メモリ ブロックに配列からの要素がフルになるまで配置され、次のメモリ ブロックに移動してそれが繰り返されます。
    • complete: 完全分割では、配列を個々の要素に分割します。1 次元配列の場合は、メモリが個々のレジスタに分割されます。デフォルトの <type>complete です。
  • <factor>: サイクリック分割では、<factor> でカーネル コードの元の配列をいくつの物理メモリに分割するかを指定します。ブロック分割では、<factor> で元の配列から各物理メモリに配置する要素の数を指定します。
    重要: 完全分割 (complete) では <factor> は指定しません。
  • <dimension>: 分割する次元を指定します。1 ~ <N> の整数を指定します。Vitis コア開発キットでは、N 次元の配列がサポートされ、配列をどの 1 つの次元でも分割できます。

例 1

次のような配列宣言があるとします。

int buffer[16];

buffer という名前の整数配列には、32 ビット幅の値が 16 個格納されています。この配列に、次の宣言を使用してサイクリック分割を適用します。

int buffer[16] __attribute__((xcl_array_partition(cyclic,4,1)));

この例では、<partition_type> を cyclic に指定しているので、Vitis で配列の内容が 4 つの物理メモリに分割されます。この属性により、配列 buffer へのアクセスに使用可能なメモリ帯域幅が 4 倍になります。

Vitis コア開発キットでは、計算ユニット内すべての配列で最高 2 つの同時アクセスを保持できます。コードの元の配列を 4 つの物理メモリに分割することにより、計算ユニットで配列 buffer に対する最高 8 つの同時アクセスを保持できます。

例 2

例 1 と同じ整数配列に、次の宣言を使用してブロック分割を適用します。

int buffer[16] __attribute__((xcl_array_partition(block,4,1)));

ブロックのサイズは 4 なので、Vitis コンパイラで 4 つの物理メモリが生成され、各メモリに配列からのデータが順に配置されます。

例 3

例 1 と同じ整数配列に、次の宣言を使用して完全分割を適用します。

int buffer[16] __attribute__((xcl_array_partition(complete, 1)));

この例では、配列がカーネルのプログラマブル ロジックの分散 RAM または 16 個のレジスタに完全に分割されます。完全分割がデフォルトなので、次の宣言を使用しても同じ結果が得られます。

int buffer[16] __attribute__((xcl_array_partition));

これによりメモリ帯域幅が最大のインプリメンテーションが作成されますが、すべてのアプリケーションに適しているわけではありません。カーネル コードによる定数またはデータ依存インデックスを介したデータへのアクセス方法によって、元のコードと同等の機能を確実にするために Vitis コンパイラで各レジスタの周辺に構築する必要のあるサポート ロジックの量が異なります。Vitis コア開発キットでの一般的なベスト プラクティス ガイドラインとして、完全分割は、少なくとも 1 つの次元が定数インデックスを介してアクセスされる配列に適しています。

関連項目

xcl_array_reshape

説明

重要: 配列変数では、使用できる属性は 1 つのみです。XCL_ARRAY_RESHAPE 属性では多次元配列がサポートされますが、1 つの属性で再形成できるのは 1 つの次元のみです。

この属性は、配列分割と垂直配列マップを組み合わせます。

XCL_ARRAY_RESHAPE 属性は、配列を複数の小型の配列に分割する XCL_ARRAY_PARTITION の効果と、ビット幅を増やして配列の要素を連結する操作を組み合わせたものです。これにより、使用されるブロック RAM の数を削減すると共に、データへの並列アクセスを実現できます。この属性では元の配列よりも要素数が少なくビット幅の広い配列が作成され、1 クロック サイクルでアクセスできるデータ量が増加します。

次のようなコードがあるとします。

void foo (...) {
int array1[N] __attribute__((xcl_array_reshape(block, 2, 1)));
int array2[N] __attribute__((xcl_array_reshape(cycle, 2, 1)));
int array3[N] __attribute__((xcl_array_reshape(complete, 1)));
...
}

ARRAY_RESHAPE 属性を使用すると、配列が次の図に示すように変換されます。

1: ARRAY_RESHAPE


構文

この属性は、配列変数の定義と一緒に指定します。

__attribute__((xcl_array_reshape(<type>,<factor>, 
<dimension>)))

説明:

  • <type>: 次のいずれかの分割タイプを指定します。
    • cyclic: サイクリック分割では、配列が計算ユニットのロジックにより同時にアクセス可能な複数の小型の物理メモリとしてインプリメントされます。各メモリに要素が 1 つずつ配置され、すべての配列に配置されたら最初の配列に戻って、配列が完全に分割されるまでそれが繰り返されます。
    • block: ブロック分割では、配列が計算ユニットのロジックにより同時にアクセス可能な複数の小型メモリとしてインプリメントされます。メモリ ブロックに配列からの要素がフルになるまで配置され、次のメモリ ブロックに移動してそれが繰り返されます。
    • complete: 完全分割では、配列を個々の要素に分割します。1 次元配列の場合は、メモリが個々のレジスタに分割されます。デフォルトの <type>complete です。
  • <factor>: サイクリック分割では、<factor> でカーネル コードの元の配列をいくつの物理メモリに分割するかを指定します。ブロック分割では、<factor> で元の配列から各物理メモリに配置する要素の数を指定します。
    重要: 完全分割 (complete) では <factor> は指定しません。
  • <dimension>: 分割する次元を指定します。1 ~ <N> の整数を指定します。Vitis コア開発キットでは、N 次元の配列がサポートされ、配列をどの 1 つの次元でも分割できます。

例 1

次の例では、17 個の要素を含む 8 ビット配列 AB[17] を、ブロック マップを使用して 5 つの要素を含む 32 ビット配列に再形成 (分割およびマップ) しています。

int AB[17] __attribute__((xcl_array_reshape(block,4,1)));
ヒント: <factor> を 4 に指定すると、配列は 4 つに分割されます。つまり、17 個の要素がビット幅が 4 倍の 5 つの要素を含む配列に再形成されます。この場合、最後の要素 AB[17] は、5 番目の要素の下位 8 ビットにマップされ、5 番目の要素の残りは空になります。

例 2

次の例では、2 次元配列 AB[6][4] を次元 [6][2] の配列 1 つに再形成しています。この次元 2 のビット幅は 2 倍です。

int AB[6][4] __attribute__((xcl_array_reshape(block,2,2)));

例 3

次の例では、関数 foo の 3 次元の 8 ビット配列 AB[4][2][2] を 128 ビット幅は (4×2×2×8) の 1 要素配列 (1 つのレジスタ) に再形成しています。

int AB[4][2][2] __attribute__((xcl_array_reshape(complete,0)));
ヒント: <dimension> を 0 に指定すると、配列のすべての次元が再形成されます。

関連項目

xcl_dataflow

説明

タスク レベルのパイプライン処理をイネーブルにして関数およびループを重複できるようにし、RTL インプリメンテーションでの同時実行性を増加してデザイン全体のスループットを向上します。

C 記述では、すべての演算が順次に実行されます。pragma HLS allocation などのリソースを制限する指示子を指定しない場合、Vivado HLS ツールではレイテンシを最小限に抑え、同時実行性を向上するように処理されます。ただし、データ依存性のためにこれが制限されることがあります。たとえば、配列にアクセスする関数またはループは、完了する前に配列への読み出し/書き込みアクセスをすべて終了する必要があります。そのため、そのデータを消費する次の関数またはループの演算を開始できません。データフロー最適化を使用すると、前の関数またはループがすべての演算を完了する前に、次の関数またはループの演算を開始できるようになります。

データフロー最適化を指定した場合、HLS ツールで順次関数またはループ間のデータフローが解析され、プロデューサー関数またはループが完了する前にコンシューマー関数またはループの演算を開始できるように、ピンポン RAM または FIFO に基づいてチャネルが作成されます。これにより関数またはループを並列実行でき、レイテンシが削減されて RTL のスループットが向上します。

開始間隔 (II) (関数またはループの開始から次の関数またはループの開始までのサイクル数) が指定されていない場合は、HLS ツールで開始間隔が最小になるようにし、データが使用可能になったらすぐに演算を開始できるようにすることが試みられます。

ヒント: HLS ツールには、データフロー コンフィギュレーション設定があります。config_dataflow コマンドは、データフロー最適化で使用されるデフォルトのメモリ チャネルと FIFO の深さを指定します。詳細は、『Vivado Design Suite ユーザー ガイド: 高位合成』 (UG902) を参照してください。

DATAFLOW 最適化が機能するようにするには、デザイン内でデータが 1 つのタスクから次のタスクに流れる必要があります。次のコーディング スタイルを使用すると、HLS ツールで DATAFLOW 最適化が実行されなくなります。

  • シングル プロデューサー コンシューマー違反
  • タスクのバイパス
  • タスク間のフィードバック
  • タスクの条件付き実行
  • 複数の exit 条件を持つループ
重要: これらのコーディング スタイルのいずれかが使用されている場合、HLS ツールでメッセージが表示され、DATAFLOW 最適化は実行されません。

詳細は、『Vivado Design Suite ユーザー ガイド: 高位合成』 (UG902) を参照してください。

最後に、DATAFLOW 最適化には階層インプリメンテーションはありません。サブ関数またはループに DATAFLOW 最適化が有益な可能性のあるタスクが含まれる場合、DATAFLOW 最適化をそのループまたはサブ関数に適用するか、サブ関数をインライン展開する必要があります。

構文

XCL_DATAFLOW 属性は、関数定義またはループ定義の前に指定します。

__attribute__((xcl_dataflow))

次の例では、関数 foo 内にデータフロー最適化を指定しています。

__attribute__((xcl_dataflow))
void foo ( a, b, c, d ) {
...
}

関連項目

xcl_latency

説明

XCL_LATENCY 属性は、関数、ループ、および領域の完了するまでの最小レイテンシまたは最大レイテンシ、あるいはその両方を指定します。レイテンシは、出力を生成するのに必要なクロック サイクル数として定義されます。関数または領域のレイテンシは、コードがすべての出力値を計算して戻るまでに必要なクロック サイクル数です。ループ レイテンシは、ループのすべての反復を実行するのにかかるサイクル数です。『Vivado Design Suite ユーザー ガイド: 高位合成』 (UG902) の「パフォーマンス メトリクスの例」を参照してください。

Vivado HLS ツールでは常に、デザインのレイテンシを最短にするよう試みられます。XCL_LATENCY 属性を指定すると、ツールで次のように処理されます。

  • レイテンシが最小値より大きく最大値未満: 制約は満たされています。これ以上の最適化は実行されません。
  • レイテンシが最小値未満: HLS ツールで指定の最小レイテンシ未満を達成できる場合は、レイテンシが指定値まで拡張されます。リソース共有が増加する可能性があります。
  • レイテンシが最大値を超える: HLS ツールで最大値以下でスケジューリングできない場合は、指定された制約を満たすことができるようエフォート レベルが上げられます。それでも最大レイテンシを満たすことができない場合は、警告が表示され、達成可能な最短のレイテンシでデザインが作成されます。
ヒント: XCL_LATENCY 属性を使用すると、ツールで最良のソリューションを探すエフォートを制限することもできます。コード内のループ、関数、または領域にレイテンシ制約を指定すると、そのスコープ内で可能なソリューションが削減され、ツールの実行時間が短縮されます。詳細は、『Vivado Design Suite ユーザー ガイド: 高位合成』 (UG902) の「ランタイムおよび容量の改善」を参照してください。

構文

XCL_LATENCY 属性は、関数、ループ、または領域本体の前に指定します。
__attribute__((xcl_latency(min, max)))

説明:

  • <min>: 関数、ループ、またはコードの領域の最小レイテンシを指定します。
  • <max>: 関数、ループ、またはコードの領域の最大レイテンシを指定します。

例 1

次の例では、関数 testfor ループ最小レイテンシを 4、最大レイテンシを 8 に指定しています。

__kernel void test(__global float *A, __global float *B, __global float *C, int id) 
{
  for (unsigned int i = 0; i < id; i++)
__attribute__((xcl_latency(4, 12))) {
   C[id] = A[id] * B[id];
 }
}

関連項目

xcl_loop_tripcount

説明

XCL_LOOP_TRIPCOUNT 属性をループに適用すると、ループで実行される反復回数の合計を手動で指定できます。

重要: XCL_LOOP_TRIPCOUNT 属性は解析専用で、合成結果には影響しません。

Vivado 高位合成 (HLS) により、各ループの合計レイテンシ、つまりループのすべての反復を実行するためのクロック サイクル数がレポートされます。ループ レイテンシは、ループ反復数 (トリップカウント) に依存します。

トリップカウントは、定数値であることもあり、ループ式 (x<y など) で使用される変数の値やループ内の制御文によって異なる場合もあります。HLS ツールでトリップカウントを決定できないこともあり、その場合はレイテンシは不明になります。これは、トリップカウントの決定に使用される変数が次のいずれかの場合です。

  • 入力引数。
  • ダイナミック演算により算出される変数。

ループのレイテンシが不明または算出できない場合、XCL_LOOP_TRIPCOUNT 属性を使用してループの反復回数の最小値、最大値、および平均値を指定できます。これにより、ループのレイテンシがデザインの総レイテンシのどの程度を占めているのかがツールで解析されてレポートされるので、デザインに適切な最適化を判断するのに役立ちます。

構文

OpenCL ソースのループ宣言の前に配置します。

__attribute__((xcl_loop_tripcount(<min>, <max>, <average>)))

説明:

  • <min>: ループの反復回数の最小値を指定します。
  • <max>: ループの反復回数の最大値を指定します。
  • <avg>: ループの反復回数の平均値を指定します。

次の例では、関数 f の WHILE ループの最小トリップカウントを 2、最大トリップカウントを 44、平均トリップカウントを 33 に指定しています。

__kernel void f(__global int *a) {
unsigned i = 0;
__attribute__((xcl_loop_tripcount(2, 64, 33)))
  while(i < 64) {
    a[i] = i;
    i++;
  }
}

関連項目

xcl_max_work_group_size

説明

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 次元行列の各次元を表します。

例 1

次は、最適化されていない加算器のカーネル ソース コード例です。このデザインでは、ワーク サイズを行列のサイズ (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) のカーネルで演算できます (各実行でデータの行全体または列全体を演算)。

関連項目

xcl_pipeline_loop

説明

ループをパイプライン処理して、レイテンシを向上し、カーネル スループットおよびパフォーマンスを最大限にします。

ループを展開することにより同時実行性は増加しますが、カーネル データパスのすべての要素を常にビジー状態に保持する問題は解決されません。ループ展開されていても、ループ制御依存により順次実行となることがあります。演算が順次実行されると、ハードウェアがアイドル状態となり、パフォーマンスが低下します。

ザイリンクスでは、この問題に対処するため、XCL_PIPELINE_LOOP 属性を使用したループ パイプライン処理用に OpenCL 2.0 API 仕様上にベンダー拡張を導入しました。

デフォルトでは、v++ コンパイラにより自動的に、トリップカウントが 64 を超える場合はループがパイプライン処理され、トリップカウントが 64 以下の場合はループが展開されます。これにより、良い結果が得られるはずです。ループの前に NOUNROLL 属性と XCL_PIPELINE_LOOP 属性を指定すると、ループをパイプライン処理するよう指定できます。

構文

OpenCL ソースのループ宣言の前に配置します。

__attribute__((xcl_pipeline_loop(<II_number>)))

説明:

  • <II_number>: パイプラインの開始間隔 (II) を指定します。Vivado HLS ツールでこの指定を満たすよう試みられますが、データ依存性によって実際の開始間隔 (II) がこれより大きくなる場合があります。II を指定しない場合、デフォルトは 1 です。

例 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;
}

関連項目

xcl_pipeline_workitems

説明

ワーク アイテムをパイプライン処理してレイテンシおよびスループットを向上します。ワーク アイテムのパイプライン処理は、カーネルのワーク グループに対するループのパイプライン処理の拡張です。これは、カーネルのスループットおよびパフォーマンスを最大にするために必要です。

構文

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];
}
}

関連項目

xcl_reqd_pipe_depth

説明

重要: パイプの宣言には、小文字と数字を使用する必要があります。また、パイプで使用される変数では printf() はサポートされません。

OpenCL フレームワーク 2.0 仕様には、パイプと呼ばれる新しいメモリ オブジェクトが導入されています。パイプには、FIFO として構成されたデータが格納されます。パイプを使用すると、データを外部メモリなしで FPGA 内の 1 つのカーネルから別のカーネルにストリーミングでき、全体的なシステム レイテンシを大幅に向上できます。

Vitis コア開発キットでは、パイプはすべてのカーネル関数の外部でスタティックに定義する必要があります。パイプの深さは、パイプ宣言内で XCL_REQD_PIPE_DEPTH 属性を使用して指定する必要があります。
pipe int p0 __attribute__((xcl_reqd_pipe_depth(512)));

パイプには、ノンブロッキング モードの標準 OpenCL read_pipe() および write_pipe() ビルトイン関数、またはブロッキング モードのザイリンクスの拡張 read_pipe_block() および write_pipe_block() 関数を使用してのみアクセス可能です。

重要: 1 つのパイプは、異なるカーネル内に 1 つのプロデューサーおよびコンシューマーのみを持つことができます。

パイプ オブジェクトはホスト CPU からはアクセスできません。パイプのステータスは、OpenCL get_pipe_num_packets() および get_pipe_max_packets() ビルトイン関数を使用してクエリできます。ビルトイン関数の詳細は、Khronos OpenCL Working Group の「The OpenCL C Specification」を参照してください。

構文

パイプ オブジェクトの宣言で設定します。

pipe int <id> __attribute__((xcl_reqd_pipe_depth(<n>)));

説明:

  • <id>: パイプの ID を指定します。小文字と数字のみで指定する必要があります。たとえば、<infifo1> のように指定します (<inFifo1> は不可)。
  • <n>: パイプの深さを指定します。有効な値は 16、32、64、128、256、512、1024、2048、4096、8192、16384、32768 です。

次はザイリンクス GitHub からの dataflow_pipes_ocl 例で、パイプを使用してブロッキング read_pipe_block() および write_pipe_block() 関数によりデータを 1 つの処理段階から次の処理段階に渡しています。

pipe int p0 __attribute__((xcl_reqd_pipe_depth(32)));
pipe int p1 __attribute__((xcl_reqd_pipe_depth(32)));
// Input Stage Kernel : Read Data from Global Memory and write into Pipe P0
kernel __attribute__ ((reqd_work_group_size(1, 1, 1)))
void input_stage(__global int *input, int size)
{
__attribute__((xcl_pipeline_loop))
mem_rd: for (int i = 0 ; i < size ; i++)
{
//blocking Write command to pipe P0
write_pipe_block(p0, &input[i]);
}
}
// Adder Stage Kernel: Read Input data from Pipe P0 and write the result
// into Pipe P1
kernel __attribute__ ((reqd_work_group_size(1, 1, 1)))
void adder_stage(int inc, int size)
{
__attribute__((xcl_pipeline_loop))
execute: for(int i = 0 ; i < size ; i++)
{
int input_data, output_data;
//blocking read command to Pipe P0
read_pipe_block(p0, &input_data);
output_data = input_data + inc;
//blocking write command to Pipe P1
write_pipe_block(p1, &output_data);
}
}
// Output Stage Kernel: Read result from Pipe P1 and write the result to 
// Global Memory
kernel __attribute__ ((reqd_work_group_size(1, 1, 1)))
void output_stage(__global int *output, int size)
{
__attribute__((xcl_pipeline_loop))
mem_wr: for (int i = 0 ; i < size ; i++)
{
//blocking read command to Pipe P1
read_pipe_block(p1, &output[i]);
}
}

関連項目

xcl_zero_global_work_offset

説明

clEnqueueNDRangeKernelglobal_work_offset を NULL またはすべて 0 に設定して使用する場合、この属性を使用すると global_work_offset を常に 0 に指定できます。

この属性は、次のようなメモリ アクセスがある場合にメモリのパフォーマンスを向上できます。

A[get_global_id(x)] = ...;
注記: REQD_WORK_GROUP_SIZE、VEC_TYPE_HINT、および XCL_ZERO_GLOBAL_WORK_OFFSET を一緒に指定すると、パフォーマンスを最大限にできます。

構文

カーネル用に指定されたカーネル定義またはプライマリ関数の前に配置します。

__kernel __attribute__((xcl_zero_global_work_offset))
void test (__global short *input, __global short *output, __constant short *constants) { }

関連項目