GPU計算モデル（SIMT・ワープ）

GPUは「揃った仕事の束」を流す装置

GPU が CPU より桁違いのスループットを出せるのは、クロックが速いからでも1コアが賢いからでもありません。同じ命令列をたどる大量のスレッドを束ね、制御回路を共有して演算器の割合を極端に高めた からです。CPU が分岐予測やアウトオブオーダー実行で1スレッドのレイテンシを削るのに対し、GPU は数千スレッドを走らせてレイテンシを別の仕事で埋めます。最適化の向きが正反対で、CPU 的な直感がそのままだと逆効果になります。土台となるメモリ階層やレイテンシ隠蔽の考え方は /hardware-components/ の各記事とも通じます。

演算の実体は SM（Streaming Multiprocessor、AMD では CU＝Compute Unit） で、GPU はこれを数十個並べた構造です。各 SM が独立にスレッドの集団を抱え、内部の ワープスケジューラ が毎サイクル発行するワープを選びます。

SIMTという実行モデル

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

1命令をワープ（32レーン）へ発行する様子:

  命令:  FMA  d, a, b        ← フェッチ・デコードは1回だけ
  レーン: T0  T1  T2 ... T31  ← 各スレッドが自分のレジスタで同時実行
          └────────────────┘
          同じ命令ポインタを共有する束

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

分岐発散（ダイバージェンス）のコスト

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

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

if (tid % 2 == 0) A();  else B();

  ステップ1: マスク=偶数レーン有効 → A() 実行（奇数は結果を捨てる）
  ステップ2: マスク=奇数レーン有効 → B() 実行（偶数は結果を捨てる）
  合流:      マスク=全レーン有効 に復帰（再収束）

重要なのは、発散が ワープ内で経路が割れたときだけ コストになる点です。32 スレッド全員が同じ側へ行くなら直列化は起きません。したがって、隣接スレッドが同じ分岐を取るようデータを並べる、あるいは分岐そのものを述語実行（条件付き代入）へ置き換えるのが定石です。なお NVIDIA の Volta 以降はレーンごとに命令ポインタを持つ独立スレッドスケジューリングを導入し、発散中のレーン間で前進保証が付きました。ただし再収束するまで各経路が直列に実行される コスト自体は変わりません。

メモリ階層 ── レジスタ・共有・グローバル

GPU の演算器は飢えやすく、性能は多くの場合メモリ側で決まります。各層は容量・スコープ・レイテンシが大きく異なり、どのデータをどこに置くかが設計の中心になります。

要点は速度差の桁の大きさです。レジスタは実質1サイクルで読めますが、グローバルメモリ（VRAM）は数百サイクルかかります。そこで頻繁に使うデータは 共有メモリ（AMD では LDS）へ一度読み込み、ブロック内のスレッドで使い回してグローバルへの往復を減らすのが常套手段です。共有メモリは複数バンクに分かれており、同一バンクへ別レーンが同時アクセスすると バンクコンフリクト で直列化するため、アクセスをバンクへ分散させる配置が要ります。一方でレジスタや共有メモリを1スレッド/1ブロックが使い込むほど、後述の占有率が下がるトレードオフを伴います。

メモリコアレッシング ── 帯域を決めるアクセス形

グローバルメモリへのアクセスは、スレッド個別ではなく ワープ単位でまとめて メモリトランザクションに変換されます。ワープ内の 32 スレッドが 連続した整列アドレス を参照していれば、ハードウェアはそれを少数の広いトランザクション（例：32 バイトのセクタを束ねた 128 バイト転送）へ 束ねられます。これが メモリコアレッシング です。

逆にアクセスが飛び飛び（ストライドが大きい、ランダム）だと、1ワープが多数の小トランザクションに分解され、キャッシュラインの一部しか使わないまま転送を捨てる形になり、実効帯域が桁で落ちます。メモリはそもそも固定幅のセクタ／バーストでしか動けないため、連続整列は1セクタを使い切り、飛び飛びは各セクタの端しか使わない——CPU のキャッシュ局所性と同じ原理を、ワープ全体で同時に満たす必要があると考えると腑に落ちます。

実務では「スレッド ID を最内次元（連続方向）へ対応づける」のが基本です。行優先の2次元配列 arr[row][col] なら、col を threadIdx.x に割り当て、隣接スレッドが隣接要素を読むようにします。行と列を取り違えると、隣接レーンが1行ぶん離れたアドレスを触ってストライドアクセスになり、帯域を無駄にします。

占有率による遅延の隠蔽

グローバルメモリの数百サイクルの遅延を、GPU は キャッシュで縮める のではなく 別の仕事で埋める ことで隠します。SM は同時に多数のワープを常駐させ、ワープスケジューラは発行のたびに「いま実行可能なワープ」を選びます。あるワープがメモリ応答待ちでストールしても、即座に別の準備済みワープへ切り替えて演算器を回し続けます。各ワープのレジスタとプログラム状態は常駐したままなので、切替コストはほぼゼロ です。CPU のコンテキストスイッチとは根本的に違います。

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

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

  遅延を隠すのに必要な同時ワープ数 ≒ メモリ遅延 / ワープあたりの発行間隔
    （待ち時間を並列な仕事で埋める、いわゆるレイテンシ隠蔽の目安）

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

• 1スレッドが使うレジスタ数が多い（レジスタ数が上限を分割する）
• 1ブロックが使う共有メモリが多い（共有メモリ容量が上限を分割する）
• ブロックあたりスレッド数やブロック数のハード上限に当たる

まとめ

• SIMT は、独立して見えるスカラスレッドを ワープ（NVIDIA=32、AMD のウェーブフロント=32/64）へ束ね、1命令を全レーンへ発行する実行モデル。制御を共有し演算器へ資源を集中できる。
• 分岐発散 はワープ内で行き先が割れたときだけ発生し、実行マスクで各経路を直列化する。隣接スレッドが同じ経路を取る配置や述語実行で回避する。
• メモリ階層は レジスタ（スレッド専有・最速）／共有メモリ（ブロック共有・オンチップ・ソフト管理）／グローバル（数百サイクル・広帯域） の三層で、頻用データを共有メモリへ載せて往復を削るのが基本。
• コアレッシング は連続整列アクセスを少数の広い転送へ束ねて実効帯域を決め、占有率 は常駐ワープで遅延を隠す余力を表す。ただし高占有率が常に最速とは限らず、律速の実測に基づいて調整する。

GPU 最適化は突き詰めると「束ねて発行し、揃えてアクセスし、待ちを別の仕事で埋める」の三点に集約されます。この計算モデルは行列積や畳み込みなど /ai/ の学習・推論を支える基盤でもあります。