ワープが発散し、関数を介して同期する必要がある可能性があることに少し混乱してい__syncthreads()ます。ブロック内のすべての要素は、SIMT 方式で同じコードを処理します。それらが同期していないのはどうしてでしょうか? スケジューラーと関係ありますか？ワープが異なれば計算時間も異なりますか? また、使用時にオーバーヘッドがあるのはなぜ__syncthreads()ですか?
ブロックに 12 の異なるワープがあり、そのうち 3 つが作業を終えたとします。そのため、アイドリングがあり、他のワープが計算時間を取得します。それとも、関数を実行するための計算時間が得られ__syncthreads()ますか?
まず、用語に注意しましょう。ワープダイバージェンスとは、コード内の制御構造（if、whileなど）により、単一のワープ内で異なる実行パスをとるスレッドを指します。質問は、ワープとワープスケジューリングに関係しています。
SIMTモデルは、すべてのスレッドがロックステップで実行されることを示唆している場合がありますが、そうではありません。まず第一に、異なるブロック内のスレッドは完全に独立しています。それらは、相互に任意の順序で実行できます。同じブロック内のスレッドに関する質問については、最初に、ブロックに最大1024（またはそれ以上）のスレッドを含めることができることを確認しましょう。ただし、今日のSM（SMまたはSMXはスレッドブロックを処理するGPU内の「エンジン」です） 1024個のcudaコアがあるため、SMがスレッドブロックのすべてのスレッドをロックステップで実行することは理論的には不可能です。単一のスレッドブロックは単一のSMで実行され、実行されないことに注意してくださいすべての（または複数の）SMで同時に。したがって、マシンに合計512以上のcudaコアがある場合でも、単一のスレッドブロックが単一のSMで実行されるため、それらすべてを使用して単一のスレッドブロックのスレッドを処理することはできません。（この理由の1つは、共有メモリなどのSM固有のリソースにスレッドブロック内のすべてのスレッドがアクセスできるようにするためです。）
では、どうなりますか？各SMにはワープスケジューラがあることがわかります。ワープは、グループ化され、スケジュールされ、実行される32個のスレッドのコレクションにすぎません。スレッドブロックに1024スレッドがある場合、ワープごとに32スレッドの32ワープがあります。たとえば、Fermiでは、SMには32個のCUDAコアがあるため、SMがロックステップでワープを実行することを考えるのは合理的です（つまり、何が起こるか、フェルミで）。ロックステップとは、（ワープの発散の場合、および命令レベルの並列性の特定の側面を無視して、ここで説明を単純にしようとしています...）ワープ内の命令は前の命令まで実行されないことを意味しますワープ内のすべてのスレッドによって実行されました。したがって、Fermi SMは、特定の瞬間にスレッドブロック内のワープの1つのみを実際に実行できます。そのスレッドブロック内の他のすべてのワープはキューに入れられ、準備が整い、待機しています。
これで、ワープの実行が何らかの理由でストールにヒットした場合、ワープスケジューラーはそのワープを自由に移動して、すぐに使用できる別のワープを取り込むことができます（この新しいワープは同じスレッドブロックからのものではない可能性がありますが、私はうまくいけば、スレッドブロックに32を超えるスレッドが含まれている場合、すべてのスレッドが実際にロックステップで実行されているわけではないことがわかります。一部のワープは他のワープよりも進んでいます。
この動作は、そうでない場合を除いて、通常は望ましいものです。条件が満たされるまで、スレッドブロック内のスレッドが特定のポイントを超えて進行しないようにする場合があります。これが__syncthreads()目的です。たとえば、グローバルから共有メモリにデータをコピーしていて、共有メモリに適切にデータが入力されるまで、スレッドブロックデータ処理を開始したくない場合があります。 __syncthreads()スレッドがバリアを超えて進行し、おそらく現在共有メモリに常駐しているデータの計算を開始する前に、すべてのスレッドがデータ要素をコピーする機会があることを確認します。
のオーバーヘッドに__syncthreads()は2つの種類があります。まず第一に、この組み込み関数に関連するマシンレベルの命令を処理するためだけに非常に小さなコストがかかります。次に、__syncthreads()通常、各ワープがバリアに到達するまで、ワープスケジューラとSMにスレッドブロック内のすべてのワープをシャッフルさせる効果があります。これが便利な場合は、すばらしいです。しかし、それが必要でない場合は、必要のないことをすることに時間を費やしています。したがって、__syncthreads()コードにたっぷりと振りかけるだけではないというアドバイスがあります。慎重に、必要な場所で使用してください。他のアルゴリズムほど使用しないアルゴリズムを作成できる場合は、そのアルゴリズムの方が優れている（高速である）可能性があります。