条件付きステートメントを避けるために、複数のOpenCLカーネルを作成する必要がありますか?
-
29-09-2019 - |
質問
OpenCLでは、複雑で実際のデータを操作する必要があるカーネルがあります。これを処理するために右のコード行を呼び出す条件付きステートメントを並べることができます。または、電話をかけて条件付きステートメントを呼び出しコードにプッシュする2つのカーネルを作成することもできます。
これは明らかに保守性にとって悪いことですが、パフォーマンスにとって重要ですか?
解決
それが1つの条件付きステートメントである場合、私の経験では、少なくともNvidiaハードウェアでは、パフォーマンスの違いは完全に無視できます。
基本的に、すべての(またはほとんどの)作業項目が同じコードパスに従っている限り、あなたは大丈夫です。撮影されたコードパスは、あなたの場合のカーネル引数に依存するため、すべての作業項目は同じパスに従います。
他のヒント
条件付きの場所に少し依存します。最初に読みやすさのコード、次にパフォーマンスを測定して問題であることがわかりました
例えば。 kernel_for_rgb_imageとkernel_for_abgr_imageは合理的な使用のように思えます。
最善の方法は、実際に2つのバリエーションをベンチマークしようとすることだと思います。場合によっては、複数の条件付きブロックがコンパイルされている場合、そのうちの1つだけが実行されたとしても、パフォーマンスが悪化する可能性があります。その理由は、GPRS(汎用レジスタ)です。コンパイラは、最悪の場合に必要なように、多くのレジスタを割り当てます。
このような解決策を提案できます。単一のカーネル関数を持っていますが、コンパイル時間条件付き:
__kernel void work()
{
#if VAR
// one code
#else
// another code
#endif
}
次に、カーネルを再コンパイルする必要があります true
/false
に設定 VAR
状態を変更するとき。明らかに、コンパイラにとっては2つのカーネルと違いはありませんが、コードの一部がそれらのカーネルで同じ場合、メンテナンスの方が良い場合があります。