36

分岐の分岐とそれを回避する方法について、インターネット上に散らばる多くの質問を見てきました。ただし、CUDA のしくみに関する記事を何十回も読んだ後でも、ほとんどの場合、分岐の分岐を回避することがどのように役立つかはわかりません。誰かが爪を伸ばして私に飛びかかる前に、私が「ほとんどの場合」と考えるものを説明させてください.

分岐分岐のほとんどのインスタンスには、多数の真に異なるコード ブロックが含まれているように私には思えます。たとえば、次のシナリオがあります。

if (A):
  foo(A)
else:
  bar(B)

この相違に遭遇した 2 つのスレッドがある場合、スレッド 1 が最初に実行され、パス A を使用します。次に、スレッド 2 がパス B を使用します。相違を取り除くために、上記のブロックを次のように変更します。

foo(A)
bar(B)

foo(A)スレッド 2 とbar(B)スレッド 1 で呼び出しても安全であると仮定すると、パフォーマンスの向上が期待できます。ただし、ここに私がそれを見る方法があります：

最初のケースでは、スレッド 1 と 2 が順次実行されます。これを 2 クロック サイクルと呼びます。

2 番目のケースでは、スレッド 1 と 2が並列に実行されfoo(A)、次に並列に実行bar(B)されます。これはまだ 2 クロック サイクルのように見えますが、違いは、前者の場合、foo(A)メモリからの読み取りが含まれる場合、そのレイテンシー中にスレッド 2 が実行を開始できるため、レイテンシーが隠蔽されることです。この場合、分岐分岐コードの方が高速です。

4

1 に答える 1

52

分岐の分岐を回避する唯一の方法は、すべてのスレッドがすべてのコードを実行できるようにすることであると想定しています（少なくとも、これはあなたが提供した例であり、あなたが作成した唯一の参照です）。

その場合、大きな違いはないことに同意します。

しかし、分岐の分岐を回避することは、if ステートメントを追加または削除して、コードをすべてのスレッドで実行するのに「安全」にするだけでなく、より高いレベルでアルゴリズムを再構築することに関係している可能性があります。

一例を挙げます。奇数のスレッドがピクセルの青のコンポーネントを処理する必要があり、偶数のスレッドが緑のコンポーネントを処理する必要があることがわかっているとします。

#define N 2 // number of pixel components
#define BLUE 0
#define GREEN 1
// pixel order: px0BL px0GR px1BL px1GR ...


if (threadIdx.x & 1)  foo(pixel(N*threadIdx.x+BLUE));
else                  bar(pixel(N*threadIdx.x+GREEN));

これは、すべての代替スレッドが指定されたパスを使用していることを意味しfooますbar。そのため、ワープの実行に 2 倍の時間がかかります。

ただし、色成分がおそらく 32 ピクセルのチャンクで連続するようにピクセル データを再配置すると、BL0 BL1 BL2 ... GR0 GR1 GR2 ...

同様のコードを書くことができます：

if (threadIdx.x & 32)  foo(pixel(threadIdx.x));
else                   bar(pixel(threadIdx.x));

まだ発散の可能性がありそうです。しかし、発散はワープ境界で発生するため、与えられたワープはifパスまたはパスのいずれかを実行するelseため、実際の発散は発生しません。

これは些細な例であり、おそらくばかげていますが、すべての分岐パスのすべてのコードを実行する必要のないワープ分岐を回避する方法がある可能性があることを示しています。

于 2013-06-20T21:45:03.730 に答える