13

CUDA を使用したプログラミングについて読んだほとんどの場所で、ワープ内のすべてのスレッドが同じことを行うことの重要性について言及されています。
私のコードでは、特定の条件を回避できない状況があります。次のようになります。

// some math code, calculating d1, d2
if (d1 < 0.5)
{
    buffer[x1] += 1;  // buffer is in the global memory
}
if (d2 < 0.5)
{
    buffer[x2] += 1;
}
// some more math code.

スレッドの中には、条件の 1 つに入るものもあれば、両方に入るものもあれば、どちらにも入らないものもあります。

条件の後にすべてのスレッドを再び「同じことをする」ようにするには、条件の後にそれらを同期する必要があり __syncthreads()ますか? それとも、これはどういうわけか自動的に起こりますか? 2 つのスレッドの 1 つが 1 つの操作の後ろにあるため、
2 つのスレッドが同じことをしていない可能性があります。それとも、ブランチの後で同じことを再び行わせるための舞台裏の努力はありますか?

4

4 に答える 4

36

ワープ内では、他のスレッドよりも「先に進む」スレッドはありません。条件付き分岐があり、それがワープ内の一部のスレッドによって取得され、他のスレッドでは取得されない場合 (ワープの「発散」とも呼ばれる)、他のスレッドは分岐が完了するまでアイドル状態になり、それらはすべて共通の命令で一緒に「収束」します。 . したがって、スレッドのワープ内同期のみが必要な場合、それは「自動的に」行われます。

しかし、異なるワープはこの方法では同期されません。そのため、アルゴリズムが特定の操作を多くのワープで完了する必要がある場合は、明示的な同期呼び出しを使用する必要があります (CUDA プログラミング ガイドのセクション 5.4 を参照してください)。

編集:いくつかのことを明確にするために、次のいくつかの段落を再編成しました。

ここには、実際には 2 つの異なる問題があります。命令の同期とメモリの可視性です。

  • __syncthreads() 命令の同期を強制し、メモリの可視性を確保しますが、ブロック間ではなく、ブロック内のみです (CUDA プログラミング ガイド、付録 B.6)。これは、共有メモリでの書き込み後読み取りには役立ちますが、グローバル メモリ アクセスの同期には適していません。

  • __threadfence()グローバルメモリの可視性を保証しますが、命令の同期は行わないため、私の経験では使用が制限されています (ただし、付録 B.5 のサンプルコードを参照してください)。

  • カーネル内ではグローバル命令同期はできません。f()任意のスレッドを呼び出す前にすべてのスレッドで実行する必要がある場合は、とを 2 つの異なるカーネルg()に分割し、それらをホストからシリアルに呼び出します。f()g()

  • 共有カウンターまたはグローバル カウンターをインクリメントするだけの場合は、アトミック インクリメント関数atomicInc()(付録 B.10) の使用を検討してください。上記のコードの場合、x1およびx2が (グリッド内のすべてのスレッドにわたって) グローバルに一意でない場合、非アトミックなインクリメントは、付録 B.2.4 の最後の段落と同様に競合状態になります。

最後に、グローバル メモリに対する操作、特に同期関数 (アトミックを含む) はパフォーマンスに悪いことに注意してください。

解決しようとしている問題を知らなければ、推測するのは困難ですが、一部の場所でグローバル メモリの代わりに共有メモリを使用するようにアルゴリズムを再設計できる可能性があります。これにより、同期の必要性が減り、パフォーマンスが向上します。

于 2009-10-29T17:07:03.860 に答える
2

CUDAベストプラクティスガイドのセクション6.1から：

フロー制御命令（if、switch、do、for、while）は、同じワープのスレッドを発散させることにより、命令スループットに大きな影響を与える可能性があります。つまり、さまざまな実行パスをたどります。これが発生した場合は、さまざまな実行パスをシリアル化する必要があり、このワープに対して実行される命令の総数が増加します。すべての異なる実行パスが完了すると、スレッドは同じ実行パスに収束します。

したがって、特別なことをする必要はありません。

于 2009-10-29T16:57:15.943 に答える
2

ガブリエルの応答では：

「グローバル命令同期はカーネル内では不可能です。任意のスレッドで g() を呼び出す前にすべてのスレッドで f() を実行する必要がある場合は、f() と g() を 2 つの異なるカーネルに分割し、それらをホストからシリアルに呼び出します。 "

同じスレッドで f() と g() が必要な理由が、レジスタ メモリを使用しているためであり、f からのレジスタまたは共有データを g に取得したい場合はどうすればよいでしょうか? つまり、私の問題では、ブロック間で同期する全体的な理由は、f からのデータが g で必要とされるためです。カーネルに分割すると、レジスタ データを f から g に転送するために大量の追加のグローバル メモリが必要になります。避けたい

于 2012-12-07T22:34:14.617 に答える
1

The answer to your question is no. You don't need to do anything special. Anyway, you can fix this, instead of your code you can do something like this:

buffer[x1] += (d1 < 0.5);
buffer[x2] += (d2 < 0.5);

You should check if you can use shared memory and access global memory in a coalesced pattern. Also be sure that you DON'T want to write to the same index in more than 1 thread.

于 2010-02-09T00:05:50.083 に答える