スレッド/ブロックのレイアウトについては、 CUDA プログラミング ガイドで詳しく説明されています。特に、第 4 章では次のように述べています。
CUDA アーキテクチャは、マルチスレッド ストリーミング マルチプロセッサ (SM) のスケーラブルなアレイを中心に構築されています。ホスト CPU 上の CUDA プログラムがカーネル グリッドを呼び出すと、グリッドのブロックが列挙され、利用可能な実行能力を持つマルチプロセッサに分散されます。スレッド ブロックのスレッドは 1 つのマルチプロセッサで同時に実行され、複数のスレッド ブロックは 1 つのマルチプロセッサで同時に実行できます。スレッド ブロックが終了すると、空のマルチプロセッサで新しいブロックが起動されます。
各 SM には 8 つの CUDA コアが含まれており、一度に 32 スレッドの 1 つのワープを実行しているため、ワープ全体に対して 1 つの命令を発行するのに 4 クロック サイクルかかります。任意のワープのスレッドがロックステップで実行されると想定できますが、ワープ間で同期するには、 を使用する必要があります__syncthreads()。
Compute Work Distributor は、SM にスレッド ブロック用の十分なリソース (共有メモリ、ワープ、レジスタ、バリアなど) がある場合にのみ、SM でスレッド ブロック (CTA) をスケジュールします。共有メモリなどのスレッド ブロック レベルのリソースが割り当てられます。割り当ては、スレッド ブロック内のすべてのスレッドに十分なワープを作成します。リソース マネージャーは、ラウンド ロビンを使用してワープを SM サブパーティションに割り当てます。各 SM サブパーティションには、ワープ スケジューラ、レジスタ ファイル、および実行ユニットが含まれます。ワープがサブパーティションに割り当てられると、それが完了するか、コンテキスト スイッチ (Pascal アーキテクチャ) によって横取りされるまで、サブパーティションに残ります。コンテキスト スイッチの復元では、ワープは同じ SM の同じワープ ID に復元されます。
ワープ内のすべてのスレッドが完了すると、ワープ スケジューラは、ワープによって発行されたすべての未処理の命令が完了するのを待ち、リソース マネージャーは、ワープ ID とレジスタ ファイルを含むワープ レベルのリソースを解放します。
スレッド ブロック内のすべてのワープが完了すると、ブロック レベルのリソースが解放され、SM は Compute Work Distributor にブロックが完了したことを通知します。
ワープがサブパーティションに割り当てられ、すべてのリソースが割り当てられると、ワープはアクティブであると見なされます。これは、ワープ スケジューラがワープの状態をアクティブに追跡していることを意味します。各サイクルで、ワープ スケジューラは、どのアクティブなワープがストールしているか、またどのワープが命令を発行する資格があるかを判断します。ワープ スケジューラは、最も優先度の高い適格なワープを選択し、ワープから 1 ～ 2 の連続した命令を発行します。二重発行のルールは、各アーキテクチャに固有です。ワープがメモリ ロードを発行する場合、依存する命令に到達するまで、独立した命令を実行し続けることができます。ロードが完了するまで、ワープはストールしていると報告します。従属演算命令についても同様です。SM アーキテクチャは、ワープ間のサイクルごとに切り替えることで、ALU とメモリのレイテンシの両方を隠すように設計されています。
この回答では、CUDA コアという用語を使用していません。これは、誤ったメンタル モデルを導入するためです。CUDA コアは、パイプライン化された単精度浮動小数点/整数実行ユニットです。発行率と依存関係の待ち時間は、各アーキテクチャに固有です。各 SM サブパーティションと SM には、ロード/ストア ユニット、倍精度浮動小数点ユニット、半精度浮動小数点ユニット、分岐ユニットなど、他の実行ユニットがあります。
パフォーマンスを最大化するために、開発者はブロック、ワープ、レジスター/スレッドのトレードオフを理解する必要があります。
占有率という用語は、SM 上の最大ワープに対するアクティブなワープの比率です。Kepler - Pascal アーキテクチャ (GP100 を除く) には、SM ごとに 4 つのワープ スケジューラがあります。SM ごとのワープの最小数は、ワープ スケジューラの数と少なくとも等しくなければなりません。アーキテクチャの依存実行レイテンシが 6 サイクル (Maxwell と Pascal) の場合、レイテンシをカバーするには、SM ごとに 24 (24 / 64 = 占有率 37.5%) であるスケジューラごとに少なくとも 6 つのワープが必要です。スレッドに命令レベルの並列性がある場合、これを減らすことができます。ほとんどすべてのカーネルは、80 ～ 1000 サイクルかかる可能性があるメモリ ロードなどの可変レイテンシ命令を発行します。これには、遅延を隠すために、ワープ スケジューラごとにより多くのアクティブなワープが必要です。カーネルごとに、ワープ数と共有メモリやレジスタなどの他のリソースとの間にトレードオフ ポイントがあるため、100% の占有率を最適化することは推奨されません。他の犠牲が生じる可能性があるためです。CUDA プロファイラーは、開発者がそのバランスを判断するのに役立つように、命令の発行率、占有率、およびストールの理由を特定するのに役立ちます。
スレッド ブロックのサイズは、パフォーマンスに影響を与える可能性があります。カーネルに大きなブロックがあり、同期バリアを使用している場合、バリア ストールがストールの原因になる可能性があります。これは、スレッド ブロックごとのワープを減らすことで軽減できます。