私の理解では、ワープは実行時にタスク スケジューラを介して定義されるスレッドのグループです。CUDA のパフォーマンスに重要な部分の 1 つは、ワープ内のスレッドの分岐です。ハードウェアがワープを構築する方法を適切に推測する方法はありますか?スレッドブロック内？
たとえば、スレッド ブロック内の 1024 スレッドでカーネルを開始しました。ワープはどのように配置されているのでしょうか。スレッド インデックスからそれを知ることができますか (または、少なくとも適切な推測を行うことができますか)。
これを行うことにより、特定のワープ内のスレッドの発散を最小限に抑えることができるためです。
ワープ内のスレッドの配置は実装に依存しますが、atm は常に同じ動作を経験しています。
ワープは 32 スレッドで構成されますが、ワープ スケジューラはワープを半分にするための命令を 1 回 (16 スレッド) 発行します。
1D ブロックを使用する場合 (threadIdx.x 次元のみが有効)、ワープ スケジューラは、threadIdx.x = (0..15) (16..31) ...に対して 1 つの命令を発行します。
2D ブロックを使用する場合 (threadIdx.x および threadIdx.y ディメンションが有効)、ワープ スケジューラは次の方法で発行を試みます。
threadIdx.y = 0 threadIdx.x = (0 ..15) (16..31) ... など
そのため、連続した threadIdx.x コンポーネントを持つスレッドは、同じ命令を 16 のグループで実行します。
ワープは、同時に実行される32のスレッドで構成されます。いつでも32のバッチがGPUで実行され、これはワープと呼ばれます。
次に実行するワープを制御できると述べているところはどこにも見つかりません。知っているのは、32スレッドで構成されており、スレッドブロックは常にその数の倍数である必要があるということだけです。
単一のブロック内のスレッドは、ソフトウェアデータキャッシュを共有する単一のマルチプロセッサで実行され、同じブロック内のスレッドとデータを同期および共有できます。ワープは常に単一のブロックからのスレッドのサブセットになります。
メモリ操作とレイテンシに関しては、これもあります。
ワープ内のスレッドがデバイスメモリ操作を発行すると、メモリレイテンシが長いため、その命令には非常に長い時間がかかり、おそらく数百クロックサイクルかかります。主流のアーキテクチャは、レイテンシを減らすためにキャッシュメモリ階層を追加します。Fermiにはいくつかのハードウェアキャッシュが含まれていますが、ほとんどのGPUは、キャッシュメモリが効果的でないストリームまたはスループットコンピューティング用に設計されています。代わりに、これらのGPUは、高度なマルチスレッドを使用することでメモリレイテンシを許容します。Teslaは各マルチプロセッサで最大32のアクティブワープをサポートし、Fermiは最大48をサポートします。メモリ操作で1つのワープが停止すると、マルチプロセッサは別のレディワープを選択してそのワープに切り替えます。このように、コアをビジー状態に保つのに十分な並列処理がある限り、コアは生産的である可能性があります。
スレッドブロックをワープに分割することに関して、私はこれを見つけました：
ブロックが2Dまたは3Dの場合、スレッドは1次元、2次元、3次元の順に並べられ、次に32のワープに分割されます。