GPUメモリ階層と実行モデル図

2つの階層を重ねて見る

GPU プログラミングが難しいのは、実行の階層とメモリの階層という独立した2軸を、頭の中で重ね合わせる必要があるからです。スレッドの束ね方（実行）と、データの置き場所（メモリ）を別々に最適化することはできません。両者はスコープという概念で結びついています。

実行階層は次の入れ子構造です。

グリッド（Grid）           ← カーネル起動1回ぶんの全スレッド
└─ ブロック（Block）       ← 1つのSM上で協調する数百スレッドの集団
   └─ ワープ（Warp, 32本） ← 1命令を全レーンに流す実行単位（SIMT）
      └─ スレッド（Thread） ← 個々のレーン

グリッドは複数のブロックから成り、各ブロックは1つの SM（Streaming Multiprocessor） に割り当てられて、そこに常駐したまま実行されます。ブロック内のスレッドはハードウェアによって自動的に 32 本ずつのワープへ刻まれ、ワープ単位で命令が発行されます。SIMT 実行そのものはGPUアーキテクチャの原理で詳しく扱っているので、ここでは「どの層がどのメモリを見るか」に絞ります。

実行スコープとメモリのマップ

各メモリ空間は、上の実行階層のどの層から共有されるかで性質が決まります。これが GPU メモリ階層の本質です。

ポイントは、可視スコープと速度が連動していることです。狭いスコープ（スレッド／ブロック）に閉じたメモリほど物理的に演算器の近くにあり、速くて狭い。広いスコープ（グリッド全体）から見えるメモリほど遠くにあり、遅くて広い。最適化とは、データをできるだけ狭いスコープに引き上げて再利用することに他なりません。

レイテンシと帯域の桁感

数字の桁を押さえると判断が速くなります。具体値は世代で変わりますが、比の構造は安定しています。

アクセスコスト（おおよその比）
レジスタ        : ~1サイクル
共有メモリ/L1   : ~20-30サイクル
L2              : ~200サイクル
グローバル(HBM) : ~400-600サイクル（DRAM行ミス時はさらに）

帯域（おおよその比）
レジスタ/共有メモリ（オンチップ集計） >> HBM帯域

グローバルメモリ（VRAM。近年のハイエンドは HBM＝広帯域メモリを積層で接続）は、TB/s 級の帯域を持つ一方でレイテンシは数百サイクルあります。GPU はこの遅延をキャッシュで縮めるより、待機ワープへの切替で別の仕事を詰めて隠す設計です（占有率による隠蔽）。とはいえ帯域は有限なので、グローバルへの往復回数そのものを減らすことが効きます。

再利用パターン：グローバル→共有のタイリング

実行階層とメモリ階層を結ぶ典型例がタイリングです。行列積を例に、擬似コードで「どの層がどのメモリに触るか」を示します。

// ブロック内の全スレッドで協調してタイルを共有メモリへ載せる
__shared__ float tileA[16][16];   // スコープ＝ブロック
__shared__ float tileB[16][16];
float acc = 0;                    // スコープ＝スレッド（レジスタ）

for (各タイル k) {
    // 1) グローバル(HBM)から共有メモリへ、コアレスして読み込む
    tileA[ty][tx] = A_global[...];   // 隣接スレッドが隣接アドレス
    tileB[ty][tx] = B_global[...];
    __syncthreads();                 // ブロック内の全ワープを同期

    // 2) 以降は高速な共有メモリから何度も読む（再利用）
    for (i in 0..15) acc += tileA[ty][i] * tileB[i][tx];
    __syncthreads();
}
C_global[...] = acc;                 // 結果を1回だけグローバルへ書く

この構造では、各グローバル要素を1回だけ読んで共有メモリに置き、ブロック内で何度も再利用します。素朴な実装が同じ行・列を繰り返しグローバルから読むのに対し、タイリングはグローバルアクセスをタイル幅ぶんの一回に圧縮します。__syncthreads() がブロックという実行スコープと共有メモリというメモリスコープを噛み合わせる関節です。ブロックを超えた同期は（原則）できないため、ブロック粒度がタイルの自然な単位になります。

容量はスコープで奪い合う

メモリ階層のもう一つの制約は容量の有限性です。SM 上のレジスタファイルと共有メモリは、そこに常駐する全ワープで分け合います。

• 1スレッドが使うレジスタ数が多い → SM に同時に載るスレッド数（＝ワープ数）が減る
• 1ブロックが使う共有メモリが多い → SM に同時に載るブロック数が減る

どちらも同時常駐ワープ数（占有率）を下げ、遅延を隠す余力を削ります。逆に共有メモリを使わなさすぎると再利用が起きずグローバル帯域で詰まります。この**「オンチップ資源を再利用に回す」対「常駐数を確保する」のトレードオフ**こそ、GPU チューニングの中心です。

まとめ

• 実行はグリッド→ブロック→ワープ→スレッドの入れ子、メモリはレジスタ→共有/L1→L2→グローバルの階層で、両者は可視スコープで対応する。
• スコープが狭いメモリほど演算器に近く速くて狭く、広いメモリほど遠く遅くて広い。最適化はデータを狭いスコープへ引き上げる作業に等しい。
• 共有メモリへのタイリングでグローバル往復を減らし、コアレスアクセスで帯域を使い切り、占有率で残る遅延を隠す。
• オンチップのレジスタ・共有メモリは全ワープで奪い合うため、再利用と常駐数のトレードオフを実測で詰める。

要するに GPU 性能は「どの層に何を置き、どの層から何回触るか」というスコープ設計の問題です。土台となるキャッシュメモリの原理と、束ね方を扱うSIMTとワープ実行を合わせて押さえると、最適化の判断が速くなります。