GPUアーキテクチャの原理 ─ SIMTとワープ実行

SIMTという実行モデル

GPU は数千本のスレッドを同時に走らせて高いスループットを出します。その心臓部が SM（Streaming Multiprocessor） で、GPU はこの演算ユニットを数十個並べた構造です。各 SM が独立にスレッドの集団を実行します。

鍵は実行の粒度です。SM は個々のスレッドをバラバラに動かすのではなく、ワープ（NVIDIA で 32 スレッド。AMD ではウェーブフロントと呼び 32 または 64）という固定本数の束にまとめ、1つの命令を束の全レーンへ一斉に流します。これが SIMT（Single Instruction, Multiple Threads） です。

ワープ（32スレッド）に対する1命令の発行
命令: FMA r0, r1, r2     ← 1つの命令
レーン:  T0  T1  T2 ... T31  ← 32スレッドが各自のレジスタで同時実行

SIMD（1命令で固定長ベクトルを処理）と似ますが、SIMT では各レーンが独立したスレッドのように見え、自分のレジスタ・自分の分岐結果・自分のアドレスを持てます。プログラマは1スレッド分のスカラコードを書き、ハードウェアが裏でレーンに割り当てます。命令フェッチとデコードはワープ単位で1回なので、制御ロジックの面積と電力を 32 レーンで共有でき、演算器にトランジスタを集中させられます。

分岐ダイバージェンス

SIMT の弱点は条件分岐です。ワープは命令を共有するので、if でレーンごとに行き先が分かれると 1命令で両方を実行できません。これを分岐ダイバージェンスと呼びます。

ハードウェアは実行マスク（どのレーンが有効か示すビット列）で対処します。まず条件が真のレーンだけを有効にして then 側を実行し、次に偽のレーンを有効にして else 側を実行します。両経路を直列に通すため、最悪では実行時間が単純に倍化します。

if (tid % 2 == 0) A();  else B();
ステップ1: マスク=偶数レーン → A() 実行（奇数は待機）
ステップ2: マスク=奇数レーン → B() 実行（偶数は待機）
合流後:    マスク=全レーン

重要なのは、ダイバージェンスはワープ内で経路が割れたときだけコストになる点です。32 スレッドが全員 then に行くなら直列化は起きません。したがって、隣接スレッドが同じ分岐を取るようデータを並べる、あるいは分岐自体を算術（条件移動・述語実行）に置き換えるのが定石です。

メモリコアレッシング

GPU の演算器は飢えやすく、性能はしばしばメモリ帯域で決まります。グローバルメモリ（VRAM）へのアクセスは個別のスレッド単位ではなく、ワープ単位でまとめてメモリトランザクションに変換されます。

このとき、ワープ内の 32 スレッドが連続した整列アドレスを参照していれば、ハードウェアはそれを少数の広いトランザクション（例：128 バイト1回）にまとめられます。これがメモリコアレッシングです。逆にアクセスが飛び飛び（ストライドが大きい、ランダム）だと、1ワープが多数の小トランザクションに分解され、転送した大半のバイトを使い捨てる形になり、実効帯域が桁で落ちます。

実務では「スレッド ID を最内次元（連続方向）に対応づける」のが基本です。たとえば行優先の2次元配列なら、array[row][col] の col を threadIdx.x に割り当て、隣接スレッドが隣接要素を読むようにします。同じデータを何度も使うなら、SM 内の共有メモリ（ソフトウェア管理の高速オンチップ領域）へ一度コアレスして読み込み、以降はそこから使うことで、グローバルメモリへの往復を減らします。

占有率による遅延の隠蔽

グローバルメモリのアクセスは数百サイクルかかります。GPU はこの遅延をキャッシュで縮めるのではなく、別の仕事で埋めることで隠します。

SM は同時に多数のワープを抱え、ハードウェアのワープスケジューラが命令発行のたびに「いま実行可能なワープ」を選びます。あるワープがメモリ応答を待ってストールしても、スケジューラは即座に別の準備済みワープへ切り替えて演算器を回し続けます。CPU のスレッド切替と違い、各ワープのレジスタは常駐したままなので、切替コストはほぼゼロです。

この「待機ワープで遅延を埋める」能力を測る指標が占有率（occupancy） です。

占有率 = SMに常駐している実効ワープ数 / SMが扱える最大ワープ数

占有率を上げるには、SM 上に多くのワープを同時に載せる必要があります。ところが SM のレジスタファイルと共有メモリは有限の容量を全ワープで分け合うため、

• 1スレッドが使うレジスタ数が多い
• 1ブロックが使う共有メモリが多い

ほど、同時に載せられるワープ数が減り、占有率が下がります。隠せる遅延が減ると演算器の遊びが増えます。

まとめ

• SM はスレッドをワープ（32本など）に束ね、1命令を全レーンへ流す SIMT で実行する。制御を共有し演算器に資源を集中できる。
• ワープ内で行き先が割れる分岐ダイバージェンスは両経路を直列化するため、隣接スレッドが同じ経路を取るよう設計する。
• メモリコアレッシングは連続整列アクセスを少数の広い転送にまとめ、実効帯域を決める。共有メモリで往復を減らせる。
• 占有率を確保して待機ワープへ切り替えることで、メモリ遅延を別ワープの計算で隠す。ただし高占有率が常に最速とは限らない。

GPU の並列性は、突き詰めれば「束ねて発行し、揃えてアクセスし、待ちを別の仕事で埋める」の3点に集約されます。土台にあるメモリ階層の原理はキャッシュメモリの原理も合わせてどうぞ。