GPU超並列 — 数千の凡人が、一人の天才に勝つとき

ここからのPart 6は、これまでと少し趣が変わる。私たちはAIへ向かう旅の、最後の登り坂に差しかかる。その第一歩が、AIブームの物理的な立役者——GPUだ。

「AIの学習にはGPUが要る」というのは、もはや常識だ。だがなぜCPUではダメなのか。同じ計算機なのに、何が違うのか。答えは、両者がまったく逆の思想で設計されているからだ。CPUが「少数の天才」なら、GPUは「数千の凡人」。この章では、なぜ凡人の大群が天才に勝つのかを解体する。

2つの正反対な設計思想 — 歴史の必然

第2章から、私たちはCPUを見てきた。CPUはレイテンシ指向（latency-oriented）——「一つの複雑な仕事を、いかに速く終わらせるか」に全てを賭けた設計だ。少数の非常に高性能なコアが、分岐予測・投機実行・巨大なキャッシュといった贅沢な仕組みで、逐次的な処理を高速にこなす。一人の天才が、複雑な問題を次々に解いていくイメージだ。

一方、GPUは元々、画面の何百万ピクセルを同時に塗る「画像処理」のために生まれた。ピクセルの色計算は、一つ一つは単純だが、膨大な数を同時に行う必要がある。ここからスループット指向（throughput-oriented）の設計が生まれた。「一つの仕事は遅くてもいい。とにかく大量の仕事を同時に片付ける」。

\text{CPU: 少数の高性能コア} \quad\text{vs}\quad \text{GPU: 数千の単純コア}

CPU（レイテンシ指向）              GPU（スループット指向）
┌───────────────┐               ┌───────────────────────┐
│ ┌───┐ ┌───┐   │               │ ▪▪▪▪▪▪▪▪▪▪▪▪▪▪▪▪ │
│ │Core│ │Core│  │ 大きく賢い     │ ▪▪▪▪▪▪▪▪▪▪▪▪▪▪▪▪ │ 小さく単純な
│ └───┘ └───┘   │ コア数個       │ ▪▪▪▪▪▪▪▪▪▪▪▪▪▪▪▪ │ コアが数千個
│ ┌─────────┐   │               │ ▪▪▪▪▪▪▪▪▪▪▪▪▪▪▪▪ │
│ │巨大キャッシュ│  │               │ ▪▪▪▪▪▪▪▪▪▪▪▪▪▪▪▪ │
│ └─────────┘   │               │ ▪▪▪▪▪▪▪▪▪▪▪▪▪▪▪▪ │
└───────────────┘               └───────────────────────┘
 逐次処理が速い                   並列処理の総量が桁違い

どちらが優れているという話ではない。問題の性質が、どちらを使うべきかを決める。逐次的で分岐の多い処理（OS、Webサーバー）はCPU。単純な計算を大量に並列で行う処理（画像、そしてAI）はGPUだ。

厨房のアナロジー — 一流シェフ vs 大量の見習い軍団

1000人分の、ゆで卵だけを作る仕事を考えよう。

一流シェフ（CPU）は、複雑なフルコースなら誰よりも速い。だが、ゆで卵1000個を一人で作るとなると、どんなに速くても1000回ゆでねばならない。天才の技術は、この単純作業ではほとんど活きない。

見習い軍団（GPU）は、一人ひとりは卵をゆでることしかできない凡人だ。だが1000人いれば、全員が同時に1個ずつゆでる。一斉に。結果、1000個が「1個ゆでる時間」でほぼ完成する。個々は遅くても、総量（スループット）では圧勝する。

私が飲食店の大量仕込みで学んだのは、まさにこれだ。同じ単純作業が大量にあるとき、一人の熟練者に集中させるより、手分けして一斉にやる方が圧倒的に速い。ただし条件がある。その作業が、互いに独立して並列にできること。卵をゆでる順番に依存関係がないからこそ、軍団が活きる。AIの計算——大量の掛け算と足し算——は、まさにこの「独立した単純作業の大群」なのだ。

SIMT — 「全員が同じ号令で動く」

GPUの数千コアを、どう統率するのか。一つ一つに別々の命令を与えるのは非効率だ。そこでGPUは SIMT（Single Instruction, Multiple Threads）——一つの命令を、多数のスレッドが一斉に、別々のデータに対して実行する——という方式を取る。

号令は一つ、「卵をゆでろ」。だが各見習いが手にする卵（データ）は別々。全員が同じ動作を、異なるデータに適用する。これがSIMTだ。

SIMT: 1つの命令が、多数のスレッドで一斉に実行される

  命令: 「C[i] = A[i] + B[i] を計算せよ」
         │
    ┌────┼────┬────┬────┬─── ... ─┐
    ▼    ▼    ▼    ▼    ▼         ▼
  thread0 t1   t2   t3   t4  ...  t1023
  C[0]  C[1] C[2] C[3] C[4] ...  C[1023]
   ↑ 全スレッドが同じ足し算を、別々の要素に対して同時実行

GPUはスレッドをワープ（warp、32スレッド単位）でまとめて実行する。この仕組みには弱点もある。ワープ内で条件分岐が食い違う（あるスレッドはif、別のスレッドはelse）と、GPUは両方の経路を順番に実行するしかなく、性能が落ちる（分岐ダイバージェンス）。だから「全員が同じことをする」処理ほどGPUは輝き、「バラバラに分岐する」処理は苦手——ここでもCPUとの棲み分けが物理から導かれる。

レイテンシ隠蔽 — 「待ち時間」を並列で埋める

GPUのもう一つの妙技が、メモリ待ちを隠すことだ。第11章で見たように、メモリアクセスは相対的に遅い。CPUは巨大なキャッシュでこれを緩和するが、GPUは別の手を使う。

あるスレッド群がメモリからデータが来るのを待っている間、GPUは即座に別のスレッド群の計算に切り替える。待っている暇があったら、他の仕事を進める。第10章のイベント駆動と同じ「待たない」思想だが、GPUはこれをハードウェアで、桁違いの規模でやる。膨大なスレッドを常に切り替え続けることで、個々のメモリ遅延を全体のスループットで塗りつぶす。

だからGPUの性能を引き出すには、十分な数のスレッド（仕事）を常に供給し続けることが鍵になる。仕事が足りないと、待ち時間を隠せず、数千のコアが遊んでしまう。

メモリ階層と、律速するもの

GPUには独自の高速メモリ（VRAM）があるが、それでも計算コアの速度には追いつかない。現代のAI計算では、しばしば計算そのものより、メモリからデータを運ぶ速度（メモリ帯域）がボトルネックになる。これをメモリバウンドという。

\text{実効性能} = \min\left(\text{計算能力（FLOPS）},\ \frac{\text{メモリ帯域}}{\text{1計算あたり必要データ量}}\right)

数千のコアがどれだけ速くても、そこへデータを運べなければ意味がない。見習いが1000人いても、卵を運ぶ人手が足りなければ全員が手持ち無沙汰になるのと同じだ。だからGPUプログラミングでは、データの再利用（一度運んだデータを高速なオンチップメモリに置き、何度も使う）が性能の鍵になる。第21章で触れるFlashAttentionのような最適化も、突き詰めればこの「メモリ移動をいかに減らすか」の工夫だ。

CUDAコードで超並列を体感する

NVIDIAの CUDA は、GPU上の並列計算を記述するための拡張だ。ベクトルの足し算 C[i] = A[i] + B[i] を、GPUで並列実行するコードを見てみよう。

/* GPU上で各スレッドが実行する関数（カーネル関数）*/
__global__ void vecAdd(float *A, float *B, float *C, int n) {
    /* 自分が何番目のスレッドかを計算 = 担当する要素の番号 */
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        C[i] = A[i] + B[i];   /* この1行を、数千スレッドが一斉に別々のiで実行 */
    }
}

int main(void) {
    int n = 1 << 20;  /* 約100万要素 */
    /* ... GPUメモリ確保、データ転送は省略 ... */

    /* 256スレッド/ブロック × 必要なブロック数 で起動 */
    int threadsPerBlock = 256;
    int blocks = (n + threadsPerBlock - 1) / threadsPerBlock;
    vecAdd<<<blocks, threadsPerBlock>>>(d_A, d_B, d_C, n);
    /* → 100万個の足し算が、ほぼ一斉に片付く */
    return 0;
}

CPUなら100万回ループする足し算が、GPUでは「1回の命令を100万スレッドに配る」形になる。blockIdx・threadIdx で「自分が何番目か」を知り、担当するデータだけを処理する——SIMTの号令が、このコードに現れている。PyTorchやTensorFlowで .to("cuda") と書くだけでGPUが使えるのは、この低レベルなCUDAカーネルを、フレームワークが自動生成してくれているからだ。

まとめ — 「おまじない」の消し方

CPUとGPUは正反対：CPUはレイテンシ指向（少数の天才）、GPUはスループット指向（数千の凡人）
問題の性質が使い分けを決める：逐次・分岐はCPU、独立した単純計算の大群はGPU
SIMTで統率：一つの命令を数千スレッドが別々のデータに一斉適用する。分岐の食い違いは苦手
待ち時間を並列で隠す：メモリ待ちの間に別スレッドへ切り替え、スループットで遅延を塗りつぶす
しばしばメモリバウンド：計算力よりデータ搬送が律速。データ再利用が性能の鍵

「なぜAIにGPUが必要なのか」「なぜ .to('cuda') で劇的に速くなるのか」——その答えは、スループットに全振りした設計思想にあった。

だが、GPUがひたすら高速化している「AIの計算」とは、具体的に何なのか。実はその正体は、単純な数——行列の掛け算——の途方もない繰り返しだ。次章では、AIのすべての土台にある数学、行列演算とベクトル空間へ踏み込む。

「天才が一人で成し遂げられることには限界がある。だが、凡庸な力を正しく束ねたとき、そこには天才を超える何かが生まれる。GPUが証明したのは、計算の話だけではないのかもしれない。」