6

ほとんどの場合、次のような CUDA または OpenCL プログラムでブランチが必要です。

for (int i=0; i<width; i++)
{
   if( i % threadIdx.x == 0)
     quantity += i*i;
}

コードは常に (または少なくともほとんどの場合) 非分岐スタイルで書き直すことができます。

for (int i=0; i<width; i++)
{
   quantity += i*i* (i % threadIdx.x != 0);
}

トレードオフは、単一のワープ スロットで実行するか、すべてのスレッドでより多くの計算を実行するかのいずれかのようです (2 番目のケースでは、合計が常に実行されますが、値がゼロになることもあります)。

分岐操作では、可能な分岐ごとに複数のワープ スロットが必要であると仮定すると、2 番目の方が一貫して最初よりも優れていることが予想されます。私の質問は次のとおりです。コンパイラが 1) を 2) に最適化するのが理にかなっている場合、または広く適用できる基準がない場合はコンパイラに頼ることができますか?

4

3 に答える 3

3

モジュロ演算はかなりコストがかかります。モジュロを追加すると、1つのスレッドだけが実行される単一の命令よりも多くの時間がかかると確信しています。ifなしの単一の分岐ステートメントは、ステートメントelseが実行されている間だけ他のスレッドをハングさせます。GPUは非常に高速なコンテキスト切り替え用に最適化されているため、そのためのコストはほとんどありません。

ただし、長い分岐ステートメントを使用しないことをお勧めします。GPUでのシリアル計算が多すぎる（つまり、1つのスレッドがすべての作業を行う）と、並列処理の利点が失われます。

于 2012-05-15T21:59:13.433 に答える
1

私の経験では、これらの種類のエッジケースを最適化するのは完全にコンパイラライター次第です。

では、1) を 2) に変えられないケースは考えられますか? ここに 1 つ: 私は、10 スレッドごとに特定の計算を実行する方が効率的なカーネルを作成しました。条件に関係なく、同じ結果が得られます。「結果がゼロになる以外はすべて実行」です。

ただし、threadId == 0 のチェックが十分に一般的なシナリオであっても、実際に最適化されているかどうかはわかりません。実装とデバイス自体 (CPU 対 GPU) に依存することに賭けます。

上記の理由だけでなく、一連のスレッドをスケジュール/開始/停止するのにかかるコストに基づいて、ワーク スケジューラの動作が異なる可能性があるため、何が最適かを実際に確認するために試してみる必要があります。それらはすべて実行されます（そしてほとんどがゼロ/アイデンティティの結果を提供します）。

お役に立てれば！

于 2012-05-15T21:21:22.630 に答える
0

CUDA についてはあまり思い出がありませんが、ループを並列化してみませんか? 計算を追加するには、アトミック操作 [1] を使用する必要があります。これがお役に立てば幸いです！そうでない場合は申し訳ありません。

  1. アトミック操作: http://supercomputingblog.com/cuda/cuda-tutorial-4-atomic-operations/
于 2012-05-15T20:55:41.353 に答える