CUDAとは何か——GPUの1万個のコアへ、仕事を割り振る仕組み

LLMがGPUを必要とする理由で、データセンター向けGPUには1万〜2万を超えるコアが載っていると書きました。ただ、部品があるだけでは何も起きません。

その1万人に「お前は何番目の計算をやれ」と割り振る仕組みが要ります。それがCUDAです。 前の2本では名前だけ挙げて素通りしたので、ここで中身を見ます。

そもそもCUDAとは

CUDAは、NVIDIAが開発した独自の技術です。 業界で共有されている標準規格でも、他社が作ったものでもありません。NVIDIA公式ドキュメントにも、こう明記されています。

“a parallel computing platform and programming model developed by NVIDIA that enables dramatic increases in computing performance by harnessing the power of the GPU.”

（NVIDIAが開発した、並列計算のプラットフォームおよびプログラミングモデルであり、GPUの力を活用することで計算性能を劇的に高める）

噛み砕くと、GPUに計算をさせるための書き方と、それを動かす一連の仕組みです。

登場したのは2006年11月。翌2007年6月に最初の正式版が公開されました。当時のプログラミングガイドのPDFがそのまま残っていて、表紙に「Version 1.0 6/23/2007」とあります。

そして押さえておきたい点が1つ。CUDAはNVIDIAのGPUでしか動きません。 この制約が、後半の「だから、乗り換えが重い」で見る話に直結します。

「順番にやれ」から「自分の担当だけやれ」へ

CUDAで書かれたプログラムは、普通のプログラムと形が違います。実物を見たほうが早いので、並べます。

意味ベクトルの記事で示した「掛けて足す」のうち、掛けるところだけを取り出します。384組の掛け算です。CPU向けに書くと、こうなります。

for (let i = 0; i < 384; i++) {
  c[i] = a[i] * b[i]   // 0番目から383番目まで、順番に
}

同じことをCUDAで書くと、こうです。

__global__ void kakeru(float *a, float *b, float *c) {
  int i = threadIdx.x;   // 自分が何番目かを聞く
  c[i] = a[i] * b[i];    // 1組だけ掛けて、終わり
}

繰り返しの指示が、どこにもありません。 代わりに「自分が何番目か」を聞いて、その1組だけ掛けています。

書き手が用意するのは、働き手ひとり分の仕事だけです。これを384人に一斉に配る指定を添えると、384組が同時に片づきます。0番の働き手は0番目を、383番の働き手は383番目を——それぞれ自分の番号しか見ていません。

「順番にやれ」から「自分の担当だけやれ」へ。 これがCUDAの書き方の中心です。GPUの働き手が1万人でも2万人でも、書く中身は変わりません。

なお、この働き手ひとり分の仕事をカーネル、働き手そのものをスレッドと呼びます。調べ物のときはこの語で引けます。

GPUとのデータのやり取りも、CUDAの重要な役割

もう一つ、CUDAが扱っている大きな仕事があります。データの運搬です。

CPUが使うメモリと、GPUが使うメモリは別物です。 計算させたい数字は、まずGPU側のメモリへコピーしないといけません。計算が終わったら、答えを持って帰ってくる。この往復もCUDAの担当です。

PyTorchを触ったことがあれば.to("cuda")という書き方を見たかもしれません。公式ドキュメントにあるとおり、これはデータをGPU側へ移す指示です。移し忘れたまま計算しようとすると止まります。

そして、この運搬が詰まると、いくらコアがあっても遊びます。 GPUが増産できない理由で書いたHBM（高帯域メモリ）が違いを作るのは、まさにここです。掛け算をする人数より、その人たちに材料を届ける太さのほうが先に足りなくなる。

CUDAを直接書くのは、ごく一部の専門家だけ

先ほど、CUDAのコード例を見ました。ですが、実際にその形でコードを書く人は、ごく一部です。

AIのソフトウェアは層になっています。

「専用ライブラリ」の行と「CUDA」の行は、別々の人たちの話ではありません。 cuBLASやcuDNNは、中身がCUDAのカーネルで書かれています。NVIDIAは、自社が開発したCUDAの上で自社のエンジニアにカーネルを書かせ、それをcuBLASやcuDNNというライブラリに仕立てて配っています。 NVIDIAの外にも、自分でCUDAのカーネルを書く専門家はいますが（表の「ごく一部の専門家」はこちらも含みます）、ほとんどの人はそこまで下りる必要がありません。

上の段にいる人が下の段を意識せずに書けるのは、NVIDIAが、その間を埋める専用ライブラリを自前で用意しているからです。数値を並べた表どうしの計算（行列計算）をまとめたcuBLAS、ニューラルネットワークの部品を集めたcuDNN——いずれもNVIDIAが書いて配っているものです。

つまりCUDAは、言語というより地盤です。誰も足元を見ないけれど、全部がその上に載っている。

だから、乗り換えが重い

GPUが世界一の商品になるまでで、乗り換えが進まない理由はソフトウェアの蓄積だと書きました。層の話をした今なら、もう少し正確に言えます。

書いたコードを別のGPU向けに直すこと自体は、できます。 AMDはCUDAのコードを自社向けに機械変換する道具を用意していて、公式の移植ガイドまで公開されています。

重いのは、その下です。cuBLASやcuDNNに相当するものが、同じ速さで、同じ答えを返すか。 さらにその上のフレームワークが同じように動くか。層ごとに調整と確認をやり直すことになります。

移せないのではなく、移した先で「前と同じ」を確かめる作業が層の数だけ積み上がる。 これが十数年ぶんの蓄積の正体です。

まとめ

• CUDAは、NVIDIAが開発した独自の技術で、GPUに計算をさせるための書き方と、それを動かす一連の仕組み。2006年11月に登場し、2007年6月に正式版
• 書き方の中心は「繰り返し」ではなく「自分が何番目かを聞いて、その分だけやる」。働き手ひとり分だけ書けば、1万人でも2万人でも同じ
• CPUのメモリとGPUのメモリは別物で、データの運搬もCUDAの担当。ここが詰まるとコアが遊ぶ
• ほとんどの人はCUDAを直接書かない。 PyTorchやcuBLAS、cuDNNといった層が上に載っていて、書き手はそちらだけを見る
• 乗り換えが重いのは変換できないからではなく、層ごとに「前と同じ動きか」を確かめ直す作業が積み上がるから