CUDA・GPGPUプログラミングモデル

CUDA は「数万スレッドを階層で束ねる」モデル

CUDA（Compute Unified Device Architecture）は、NVIDIA GPU を汎用計算（GPGPU）に使うためのプログラミングモデルです。CPU が数十スレッドを賢く走らせるのに対し、GPU は 数万〜数十万の軽量スレッドを一斉に走らせて処理を隠蔽 します。個々のスレッドは遅くても、圧倒的な数と高いメモリ帯域で総スループットを稼ぐ——これが GPGPU の本質です。GPU 内部のハードウェア実行の仕組みは /hardware-components/ を、行列演算への応用は /ai/ を前提知識として補完できます。

プログラムは ホスト（CPU） と デバイス（GPU） の2つの世界に分かれます。ホストは制御役で、データ転送とカーネル起動を指揮する。デバイスは計算役で、__global__ 修飾した関数（カーネル）を大量のスレッドで実行します。この分業を理解することが、CUDA を正しく書く第一歩です。

グリッド・ブロック・スレッドの3階層

カーネルを起動すると、スレッドは3階層の格子として展開されます。上から グリッド → スレッドブロック → スレッド です。この階層は論理的な設計で、ハードウェアの物理構造（SM、ワープ）に写像されて実行されます。

各スレッドは自分の座標を組み込み変数から読み、担当するデータ要素を特定します。blockIdx（グリッド内のブロック位置）、blockDim（ブロックの大きさ）、threadIdx（ブロック内のスレッド位置）を組み合わせるのが定石です。

// 1次元のグローバルインデックス算出（最頻出のイディオム）
__global__ void add(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {            // グリッドが要素数より大きい場合の番兵
        c[i] = a[i] + b[i];
    }
}

if (i < n) の番兵が重要です。ブロックサイズは 32 の倍数（多くは 128 や 256）に固定するのが定石なので、要素数 n がその倍数でないとグリッドは n を超えるスレッドを含みます。境界チェックを怠ると範囲外アクセスになります。

ワープと SIMT ── 32スレッドがロックステップで動く

論理的にはスレッドは独立ですが、ハードウェアは 32 スレッドを ワープ という束にまとめ、同一命令を一斉実行します。これが SIMT（Single Instruction, Multiple Threads）です。SIMD と違い各スレッドは自分のレジスタと分岐先を持てますが、同じワープ内でスレッドが別々の分岐に進むと厄介です。

ワープダイバージェンス（分岐発散）:

  if (threadIdx.x % 2 == 0)  A();   // 偶数レーン
  else                       B();   // 奇数レーン

  ワープ内で条件が割れると A と B を「順番に」実行し、
  片方の実行中はもう片方のレーンをマスク（休止）する。
  → 実効スループットが最悪 1/2 に低下する。

つまり 分岐はワープ単位で揃える のが鉄則です。同じ 32 スレッドが同じ道を通るようにデータやインデックスを設計すれば、発散のペナルティを避けられます。データ依存の分岐をワープ境界に合わせるだけで性能が変わります。

共有メモリと __syncthreads() ── ブロック内の協調

各 SM には 共有メモリ（shared memory）という小容量（数十〜百数十 KB 級）の高速オンチップメモリがあり、__shared__ で宣言します。グローバルメモリより桁違いに低レイテンシで、同じブロック内のスレッドがデータを共有 できます。何度も再利用するデータを共有メモリに載せ、グローバルメモリへのアクセス回数を減らすのが高速化の王道です。

ただし複数スレッドが共有メモリを読み書きするとレースが起きます。そこで __syncthreads() でブロック内の全スレッドをバリア同期し、「全員が書き終えてから、全員が読む」順序を保証します。

// タイル化行列積の中核（共有メモリで再利用性を稼ぐ）
__shared__ float As[TILE][TILE];
__shared__ float Bs[TILE][TILE];

As[ty][tx] = A[/* このタイルの担当要素 */];
Bs[ty][tx] = B[/* このタイルの担当要素 */];
__syncthreads();               // タイル全体が埋まるのを待つ

float acc = 0.0f;
for (int k = 0; k < TILE; ++k) // 共有メモリ上で内積を積む
    acc += As[ty][k] * Bs[k][tx];
__syncthreads();               // 次タイル読み込み前に全員完了を待つ

共有メモリは物理的に複数の バンク に分割されています。同一ワープ内の異なるスレッドが同じバンクの別アドレスを同時に触ると バンクコンフリクト が起き、アクセスが直列化されて遅くなります。連続する 4 バイト語が連続バンクに割り当たるため、ストライドを工夫（例: 配列幅を 1 増やしてパディング）して衝突を避けます。

ホスト・デバイス間転送とメモリ階層

GPU は独立したデバイスメモリ（VRAM）を持ち、CPU のメインメモリとは物理的に別空間です。計算前に入力をデバイスへ送り、計算後に結果を戻す必要があります。転送は PCIe（や NVLink）を通り、GPU 内部の帯域に比べて桁違いに遅い ため、ここがしばしばボトルネックになります。

古典的な CUDA プログラムの流れ:

  cudaMalloc()      デバイス側メモリを確保
  cudaMemcpy(H→D)   入力をホスト→デバイスへ転送 ← 遅い経路
  kernel<<<grid, block>>>(...)   カーネル起動（非同期）
  cudaMemcpy(D→H)   結果をデバイス→ホストへ転送 ← 遅い経路
  cudaFree()        デバイスメモリ解放

kernel<<<grid, block>>>(...) の三重山括弧が 実行構成（execution configuration） で、グリッド次元とブロック次元を指定します。カーネル起動はホストに対して 非同期 で、CPU は起動命令を投げるとすぐ次に進みます。だからこそ、転送と計算を別ストリームで重ねる（オーバーラップ）余地が生まれます。

メモリは階層になっています。速い順に レジスタ（スレッド固有）→ 共有メモリ / L1（ブロック内共有）→ L2（GPU全体）→ グローバルメモリ（VRAM） です。上位ほど小容量・低レイテンシ。データの再利用性を上位階層に閉じ込めるほど速くなります。

メモリコアレッシング ── 性能を左右する最重要ポイント

グローバルメモリアクセスで最も効くのが メモリコアレッシング（coalescing）です。GPU はグローバルメモリを 32 バイトや 128 バイト単位のトランザクションで読み書きします。ワープ内 32 スレッドが 連続したアドレス を触れば、少数のトランザクションにまとまり帯域をフルに使えます。逆にバラバラのアドレスだと、1 スレッドごとに別トランザクションが発行され、実効帯域が数分の一〜十数分の一に落ちます。

コアレッシングの良否（スレッド i がアクセスするアドレス）:

  良い: a[i]              連続 → 隣接スレッドが隣接アドレス
        隣り合う32スレッドが1本の128Bトランザクションに収まる

  悪い: a[i * stride]     飛び飛び → 各スレッドが別トランザクション
        a[perm[i]]        ランダム → まとめられず帯域を浪費

実務では データ構造をアクセスパターンに合わせる ことが鍵になります。典型例が AoS（構造体の配列）と SoA（配列の構造体）の選択です。粒子ごとに {x, y, z} を持つ構造体を並べた AoS で全粒子の x だけ走査すると、x が構造体サイズ間隔で飛び飛びになりコアレッシングが崩れます。座標軸ごとに配列を分けた SoA（x[], y[], z[]）にすれば x[i] が連続し、きれいにコアレッシングします。

性能の勘所 ── オキュパンシと律速の見極め

GPU で速さを出すには、演算器を遊ばせないだけの並列度（オキュパンシ: 同時に走らせられるワープ数の充足率）を確保しつつ、メモリと演算のどちらが律速かを見極めます。

• オキュパンシ: スレッドあたりのレジスタ数やブロックあたりの共有メモリ使用量が多いと、SM に同時常駐できるブロック数が減り、メモリ待ちを隠すワープが足りなくなる。資源の使いすぎは並列度を削る。
• メモリ律速 vs 演算律速: 到達帯域で決まるのか演算スループットで決まるのかを見極める。判断枠組みは /hardware-components/ のルーフラインモデルが有効。
• レイテンシ隠蔽: GPU はスレッドを大量に走らせ、あるワープがメモリ待ちの間に別ワープを実行して待ち時間を隠す。十分なワープ数がなければ隠蔽が効かず演算器が空転する。

これらは互いに綱引きの関係にあり、共有メモリを増やせば再利用は上がるがオキュパンシは下がる、といったトレードオフを実測（プロファイラ）で詰めるのが現実的な最適化です。

まとめ

• CUDA は グリッド → ブロック → スレッドの3階層 でカーネルを起動し、実行時は 32スレッドのワープ単位で SIMT として走る。
• ブロックは独立（相互同期しない）ことで SM 数に依らずスケールし、ワープ内の分岐発散 はスループットを削るので分岐は 32 スレッドで揃える。
• 共有メモリはブロック内共有の高速スクラッチパッドで、__syncthreads() で順序を保証し、バンクコンフリクトを避けて使う。
• ホスト・デバイス間の PCIe 転送は最大級のボトルネックで、pinned メモリ・ストリーム・往復削減で隠す。
• メモリコアレッシング（隣接スレッド＝連続アドレス、SoA 化）が性能の要で、外すと実効帯域が激減する。オキュパンシとレイテンシ隠蔽を保ちつつ、メモリ律速か演算律速かを見極めて最適化する。