私はGeForceGTX460SEを持っているので、6 SM x48CUDAコア=288CUDAコアです。1つのワープには32のスレッドが含まれ、1つのブロックで同時に（一度に）実行できるのは1つのワープのみであることが知られています。つまり、単一のマルチプロセッサ（SM）では、48コアが使用可能であっても、1つのブロック、1つのワープ、および32のスレッドのみを同時に実行できますか？
さらに、具体的なスレッドとブロックを配布する例は、threadIdx.xとblockIdx.xを使用できます。それらを割り当てるには、カーネル<<<ブロック、スレッド>>>（）を使用します。しかし、特定の数のWarpを割り当てて配布する方法と、それが不可能な場合は、なぜWarpについてわざわざ知る必要があるのでしょうか。
状況は、あなたが説明したものよりもかなり複雑です。
ALU (コア)、ロード/ストア (LD/ST) ユニット、および特殊機能ユニット (SFU) (画像の緑色) は、パイプライン化されたユニットです。それらは、完了のさまざまな段階で、同時に多くの計算または操作の結果を保持します。したがって、1 サイクルで、新しい操作を受け入れ、かなり前に開始された別の操作の結果を提供できます (私の記憶が正しければ、ALU では約 20 サイクル)。したがって、理論上、単一の SM には、48 * 20 サイクル = 960 ALU 操作を同時に処理するためのリソースがあり、ワープあたり 960 / 32 スレッド = 30 ワープとなります。さらに、LD/ST 操作と SFU 操作を、そのレイテンシーとスループットに関係なく処理できます。
ワープ スケジューラ (画像の黄色) は、ワープごとに 2 * 32 スレッド = サイクルごとにパイプラインに 64 スレッドをスケジュールできます。つまり、1 クロックあたりに取得できる結果の数です。そのため、48 コア、16 LD/ST、8 SFU というコンピューティング リソースが混在し、それぞれに異なるレイテンシがある場合、ワープの混在が同時に処理されます。任意のサイクルで、ワープ スケジューラは、SM の使用率を最大化するために、スケジュールする 2 つのワープを「ペア」にしようとします。
命令が独立している場合、ワープ スケジューラは、異なるブロックから、または同じブロック内の異なる場所からワープを発行できます。そのため、複数のブロックからのワープを同時に処理できます。
複雑さに加えて、リソースが 32 未満の命令を実行するワープは、すべてのスレッドを処理するために複数回発行する必要があります。たとえば、8 つの SFU があるため、SFU を必要とする命令を含むワープを 4 回スケジュールする必要があります。
この説明は簡略化されています。GPU が作業をスケジュールする方法を決定する他の制限もあります。Web で「フェルミ アーキテクチャ」を検索すると、詳細な情報を見つけることができます。
それで、あなたの実際の質問に来て、
なぜワープについて知りたがるのですか？
ワープ内のスレッド数を把握し、それを考慮することは、アルゴリズムのパフォーマンスを最大化しようとするときに重要になります。これらのルールに従わないと、パフォーマンスが低下します。
カーネル呼び出しで<<<Blocks, Threads>>>は、ワープ内のスレッド数で均等に分割されるスレッド数を選択してみてください。そうしないと、非アクティブなスレッドを含むブロックを起動することになります。
カーネルで、ワープ内の各スレッドが同じコード パスをたどるようにしてください。そうしないと、ワープ発散と呼ばれるものが発生します。これは、GPU が分岐コード パスのそれぞれを介してワープ全体を実行する必要があるために発生します。
カーネルでは、ワープ内の各スレッドが特定のパターンでデータをロードして保存するようにしてください。たとえば、ワープ内のスレッドがグローバル メモリ内の連続する 32 ビット ワードにアクセスするようにします。
スレッドは、必ず 1 ～ 32、33 ～ 64 の順序でワープにグループ化されますか?
はい、プログラミング モデルは、スレッドがその特定の順序でワープにグループ化されることを保証します。
分岐コード パスの最適化の簡単な例として、ブロック内のすべてのスレッドを 32 スレッドのグループに分離することを使用できますか? 例: switch (threadIdx.s/32) { case 0: /* 1 warp*/ break; ケース 1: /* 2 ワープ*/ ブレーク; /* その他 */ }
丁度 ：）
1 つのワープで一度に読み取らなければならないバイト数: 4 バイト * 32 スレッド、8 バイト * 32 スレッド、または 16 バイト * 32 スレッド? 私が知る限り、一度にグローバル メモリへの 1 つのトランザクションは 128 バイトを受け取ります。
はい、グローバル メモリへのトランザクションは 128 バイトです。そのため、各スレッドが連続するアドレスから 32 ビット ワードを読み取る場合 (おそらく 128 バイトにアラインする必要もあります)、ワープ内のすべてのスレッドを 1 つのトランザクションで処理できます (4 バイト * 32 スレッド = 128 バイト)。 ）。各スレッドがより多くのバイトを読み取る場合、またはアドレスが連続していない場合は、より多くのトランザクションを発行する必要があります (タッチされる 128 バイト行ごとに個別のトランザクションを使用して)。
これについては、CUDA プログラミング マニュアル 4.2、セクション F.4.2、「グローバル メモリ」で説明されています。また、L2 キャッシュには 32 バイトのキャッシュ ラインがあるため、L2 にのみキャッシュされるデータでは状況が異なるという宣伝文句もあります。データが L2 のみにキャッシュされるように調整する方法や、最終的にいくつのトランザクションが発生するかはわかりません。