パフォーマンスに悪影響を及ぼす可能性があるため、CUDA での分岐は推奨されないことを理解しています。私の仕事では、数十以上のケースを含む大きな switch ステートメントを実装しなければならないことに気づきました。
これがパフォーマンスにどれほど悪影響を与えるか、誰にもわかりません。(公式ドキュメントはあまり具体的ではありません)また、この部分をより効率的に処理する方法はありますか?
パフォーマンスに悪影響を及ぼす可能性があるため、CUDA での分岐は推奨されないことを理解しています。私の仕事では、数十以上のケースを含む大きな switch ステートメントを実装しなければならないことに気づきました。
これがパフォーマンスにどれほど悪影響を与えるか、誰にもわかりません。(公式ドキュメントはあまり具体的ではありません)また、この部分をより効率的に処理する方法はありますか?
GPU は、ワープと呼ばれる 32 のグループでスレッドを実行します。ワープ内の異なるスレッドがコード内の異なるパスを通過するたびに、GPU はコード パスごとに 1 回ずつ、ワープ全体を複数回実行する必要があります。
ワープ ダイバージェンスと呼ばれるこの問題に対処するには、特定のワープ内のスレッドが通過するコード パスの数ができるだけ少なくなるように、スレッドを調整する必要があります。それができたら、残りのワープ ダイバージェンスによるパフォーマンスの低下を受け入れるだけで済みます。場合によっては、スレッドを整理するためにできることは何もないかもしれません。その場合、異なるコード パスがカーネルまたは全体的なワークロードの大部分を占める場合、そのタスクは GPU に適していない可能性があります。
さまざまなコード パスをどのように実装するかは問題ではありません。if-else
、switch
、述語 (PTX または SASS)、ブランチ テーブル、またはその他のもの - さまざまなパスで実行されているワープ内のスレッドになると、パフォーマンスが低下します。
また、ワープ内の異なるパスの総数だけで、各パスを通過するスレッドの数も重要ではありません。
これについては、もう少し詳しく説明する別の回答があります。
複数のスイッチを回避する良い方法は、関数テーブルを実装し、スイッチの条件に基づいてインデックスごとにテーブルから関数を選択することです。__device__
CUDAを使用すると、カーネル内の関数で関数ポインターを使用できます。