CUDA Programming Guide Part 1を小学生にもわかるように読む

元記事はこちらです。

• CUDA Programming Guide Part 1 - Zenn
• 参考: NVIDIA CUDA Programming Guide
• 参考: NVIDIA Ampere Architecture In-Depth

この記事は、元記事の要約というより、「CUDAって結局なにをしているの？」を最初の一段目から理解するための補助線です。

会計士・税理士の実務感覚で言うと、CUDAは「巨大な明細表を、ものすごい人数の補助者に一斉処理してもらうための仕事の振り方」です。

最初の一言

CUDAを一言でいうと、こうです。

CPUが司令塔になり、GPUに大量の同じ作業を一気にやらせるための約束事。

たとえば、10万行の売上明細があり、各行について「税込金額を計算する」だけなら、1人のベテランが上から順に処理するより、たくさんの人に1行ずつ配ったほうが速いです。

GPUは、この「たくさんの人に同じ作業を配る」が得意です。

ここで大事なのは、CPUとGPUは同じ仕事仲間ではありますが、役割が違うことです。

なぜGPUは速いのか

CPUは「少数精鋭」です。

複雑な判断、分岐、OSやアプリ全体の管理が得意です。税務相談でいうと、論点を見つけて判断するベテランです。

GPUは「大量の単純作業チーム」です。

1人ひとりはCPUほど器用ではありません。しかし、同じ作業を同時にものすごい人数で進めます。明細10万行に対して、同じ計算式を一斉に当てるような仕事で強いです。

AIの学習や推論では、行列計算、ベクトル計算、確率計算のような「似た形の計算」が大量に出ます。だからGPUが効きます。

CUDAは言語だけではない

元記事で最初に出てくる重要ポイントは、CUDAという言葉が複数の意味で使われることです。

この記事でまず押さえるべきなのは、CUDA C++の細かい文法より前にある、CUDA Programming Modelです。

つまり、「GPUという巨大な作業会場に、どんな単位で仕事を配るのか」です。

GPUの中身を会社組織で考える

GPUの中には、SM（Streaming Multiprocessor）という小さな工場・作業島のような単位がたくさんあります。

ものすごく単純化すると、GPUはこう見れば十分です。

GPU = SMがたくさん集まったもの。

SMは、実際に計算を進める現場です。GPU全体に仕事を投げると、仕事のかたまりが各SMに配られます。

会計実務でたとえるなら、1つのGPUは大きなBPOセンターで、SMはその中の各チームです。Aチーム、Bチーム、Cチームが同時に明細を処理します。

Thread、Block、Gridを一気に理解する

CUDAでいちばん大事な3語が出てきます。

• Thread
• Thread Block
• Grid

いきなり英語で覚えようとするとつらいので、まずはこうです。

たとえば、1024行の明細を処理するとします。

1つの班に256人いるなら、4班あれば1024人です。

4班 x 256人 = 1024人

CUDA C++では、これを次のような形で指定します。

vecAdd<<<4, 256>>>(A, B, C);

これは、ざっくり言うとこうです。

4個のBlockを作って、それぞれのBlockに256個のThreadを入れて、vecAddを実行して

ここでの vecAdd は、GPU側で動く作業手順書です。

Blockの中のThreadは同じSMで働く

元記事で重要なのは、同じThread BlockにいるThreadは、同じSMで実行されるという点です。

これは、同じ班のメンバーは同じ部屋で作業する、ということです。

同じ部屋にいれば、情報共有が速いです。隣の人に「この中間計算ちょっと使わせて」と言いやすい。CUDAでは、この近い場所での共有に Shared Memory が効いてきます。

逆に、別のBlockは別のSMに割り当てられるかもしれません。しかも、どのBlockが先に走るかは基本的に保証されません。

だから、普通のCUDAプログラムでは、別Blockの結果を待つ前提で書いてはいけないと考えます。

これは、全国に散った支店で同時に作業しているようなものです。東京支店の処理が必ず大阪支店より先に終わる、という前提では設計できません。

Warpは32人1組の横並びチーム

Thread Blockの中に、さらに重要な単位があります。

Warp = 32個のThreadのまとまり。

GPUは、この32人1組に同じ命令を出します。

たとえば、32人に「自分の担当行のA列とB列を足してC列に書いて」と一斉に言うイメージです。

ここでややこしいのが、分岐です。

偶数行ならAの処理
奇数行ならBの処理

こうなると、同じ32人の中で、ある人はA処理、ある人はB処理になります。GPUは一斉号令が得意なので、こういうバラバラな動きは苦手です。

元記事で出てくる SIMT は、ここに関係します。

だから、Thread数を32の倍数にする例がよく出ます。32人1組で動くなら、256人は8 Warpぴったりです。

256人 ÷ 32人 = 8組

端数が出ると、最後の組に空席が出ます。

GPUメモリは「机、班の共有机、倉庫」で考える

GPUを理解するうえで、メモリはとても重要です。

元記事では、Global Memory、Register、Shared Memory、L1 cacheなどが出てきます。最初は全部覚えなくて大丈夫です。

まずは3段階で考えます。

PyTorchでよく見る CUDA out of memory は、だいたいこの大きな倉庫、つまりGlobal Memoryが足りない話です。

ただし、CUDAを低いレイヤーで書くと、Global Memoryだけでなく、RegisterやShared Memoryの使い方が速度に大きく効きます。

FlashAttentionのような高速化技術がすごい理由も、雑に言えば「倉庫に何度も取りに行かず、近い場所のメモリをうまく使う」からです。

KernelはGPUに配る作業手順書

CUDA C++では、GPUで動かす関数に __global__ を付けます。

__global__ void addOne(float* values) {
  int i = threadIdx.x;
  values[i] = values[i] + 1.0f;
}

これは「この関数はGPUで実行する作業手順書ですよ」という印です。

そしてCPU側から、次のように呼び出します。

addOne<<<1, 256>>>(values);

<<<1, 256>>> は、普通のC++にはないCUDA C++の書き方です。

意味はこうです。

1班、256人で、この作業を始めて

みんな同じコードを動かすのに、なぜ別々の行を処理できるのか

ここがCUDAの最初の壁です。

GPUでは、たくさんのThreadが同じKernelを実行します。

では、全員が同じことをしたら、同じ行ばかり計算してしまうのではないか。

そこで出てくるのが、threadIdx、blockIdx、blockDim です。

これらは、各Threadに配られる座席番号のようなものです。

全体で見た自分の担当番号は、次の式で出せます。

int i = blockIdx.x * blockDim.x + threadIdx.x;

会計の明細処理で言うと、これは「自分が何行目を担当するか」を計算しているだけです。

たとえば、1班256人なら、

こうやって、同じコードを動かしているのに、各Threadが別々の明細行を処理できます。

Bounds Checkingは「存在しない行を触らない」ため

元記事のベクトル加算では、次のようなチェックが出ます。

if (i < vectorLength) {
  C[i] = A[i] + B[i];
}

これは、とても実務的です。

たとえば明細が1000行しかないのに、処理人数を1024人で用意したとします。24人分は担当行がありません。

その24人が存在しない1001行目以降を読みに行くと、事故になります。

だから、

自分の担当番号が実在する行番号なら処理する
実在しないなら何もしない

というチェックを入れます。

CUDAに限らず、配列処理ではこの発想がとても大事です。

Unified MemoryとExplicit Memory Management

GPUを使うときは、CPU側のデータをGPU側から読めるようにする必要があります。

ここで2つの考え方が出ます。

Unified Memoryは楽です。

ただし、楽なぶん、いつどこでデータが動くかを細かく制御しにくいです。

Explicit Memory Managementは面倒です。

ただし、面倒なぶん、データ移動と計算を重ねたり、余計な移動を減らしたりしやすいです。高速化を詰めるなら、こちらの理解が必要になります。

CPUはGPUの完了を待たない

Kernel Launchは、基本的に非同期です。

つまり、CPUがGPUに「これやって」と指示を出したあと、GPUの作業完了を待たずにCPU側の次の処理へ進むことがあります。

会計実務で言うと、補助者チームに明細チェックを依頼した瞬間、レビュー担当者が「もう終わったはず」と思って結果表を見に行くと、まだ作業中かもしれません。

だから、必要なところで待ちます。

cudaDeviceSynchronize();

これは、かなり噛み砕くと、

GPU側に頼んだ作業が終わるまで、CPU側はここで待つ

という命令です。

会計士・税理士向けの読み替え

CUDAの用語は難しいですが、会計実務に置き換えるとかなり見通しがよくなります。

これだけ持って元記事を読み直すと、専門用語の圧がかなり下がるはずです。

最低限の全体像

最後に、CUDA Part 1の内容を1枚の流れにします。

1. CPUがデータを用意する。
2. GPUが読める場所にデータを置く。
3. CPUがKernelをLaunchする。
4. Grid、Block、Threadに仕事が分配される。
5. 各Threadが自分の担当番号を計算する。
6. 各Threadが担当データを処理する。
7. 必要ならCPUがGPUの完了を待つ。
8. 結果をCPU側で確認する。

CPU: データ準備
  ↓
GPU: 大量のThreadで一斉処理
  ↓
CPU: 結果確認

ここまでわかれば、Part 1の大枠はつかめています。

この記事で覚えることは5つだけ

最初から全部覚えようとしなくて大丈夫です。

まずは、この5つだけで十分です。

元記事を読むときの順番

元記事をいきなり上から精読すると、用語が多くて苦しくなります。

おすすめは、この順番です。

1. Thread / Block / Gridの図を見る。
2. <<<4, 256>>> の意味だけ理解する。
3. blockIdx.x * blockDim.x + threadIdx.x の式を理解する。
4. Register / Shared Memory / Global Memoryの違いを見る。
5. Unified MemoryとExplicit Memory Managementの違いを見る。
6. 最後にHardware Modelの細かい名前へ戻る。

最初からGPC、SM、Unified Data Cache、Compute Capabilityを全部覚えようとすると、たぶん止まります。

先に「仕事の配り方」を理解してから、あとでハードウェアの名前に戻るほうが楽です。

まとめ

CUDAは、最初に出てくる単語が強そうなので難しく見えます。

でも、芯はかなりシンプルです。

大量の同じような作業を、小さな作業者に分けて、GPU上で一気に処理する。

そのために、

• 仕事を配る単位として、Thread / Block / Gridがある
• 32人単位のWarpがある
• 近いメモリと遠いメモリがある
• CPUはGPUに作業を依頼し、必要なら完了を待つ

この絵を頭に置いておけば、元記事の専門用語は「知らない外国語」ではなく、「すでに知っている仕事の仕組みに付いた名前」として読めます。