3.8. 高価な機能の回避
一部の機能はFPGAで実装するのに費用がかかります。高価な機能は、カーネルのパフォーマンスを低下させるか、実装するために大量のハードウェアを必要とする可能性があります。
以下の機能は高価です。
- 整数除算とモジュロ（剰余）演算子
- 加算、乗算、絶対値、および比較を除くほとんどの浮動小数点演算子
注: 浮動小数点演算の最適化の詳細については、「浮動小数点演算の最適化のセクションを参照してください。
- アトミック関数
対照的に、安価な機能はカーネルの性能にはほとんど影響を与えず、その実装は最小のハードウェアしか消費しません。
以下の機能は安価です。
- AND、NAND、OR、NOR、XOR、およびXNORなどの2進論理演算
- 1つの定数引数による論理演算
- 定数でシフト
- 整数の乗算と2の累乗である定数による除算
高価な関数がワークグループ内のすべてのWork-Itemに対して新しいデータを生成する場合、それをカーネルでコード化することは有益です。これに対して、以下のコード例は、NDRangeの各Work-Itemで実行される高価な浮動小数点演算（除算）のケースを示しています。
__kernel void myKernel (__global const float * restrict a,
__global float * restrict b,
const float c, const float d)
{
size_t gid = get_global_id(0);
//inefficient since each work-item must calculate c divided by d
b[gid] = a[gid] * (c / d);
}
この計算の結果は常に同じです。このような冗長かつハードウェアのリソース集約的な動作を回避するには、ホスト・アプリケーションで計算を実行し、使用するNDRange内のすべてのWork-Itemの引数として結果をカーネルに渡します。変更されたコードを以下に示します。
__kernel void myKernel (__global const float * restrict a,
__global float * restrict b,
const float c_divided_by_d)
{
size_t gid = get_global_id(0);
/*host calculates c divided by d once and passes it into
kernel to avoid redundant expensive calculations*/
b[gid] = a[gid] * c_divided_by_d;
}Intel® FPGA SDK for OpenCL™オフライン・コンパイラーは、NDRange全体でWork-Itemに依存しない動作を1つの動作に統合します。次に、すべてのWork-Itemにわたって結果を共有します。最初のコード例では、オフライン・コンパイラーは、 cによるdの除算がすべてのWork-Itemにわたって一定であるため、すべてのWork-Itemで共有される1つの除算ブロックを作成します。この最適化は冗長ハードウェアの量を最小限に抑えるのに役立ちますしかし、整数除算の実装には、相当量のハードウェア・リソースが必要です。したがって、除算演算をホスト・プロセッサーにオフロードし、結果を引数としてカーネルに渡してハードウェア・リソースを節約することが有益です。
関連情報