「GPUで計算する」と聞くと、なんだか一瞬で魔法みたいに処理が終わる印象があります。
でも元記事は、その裏側にある泥臭い手順を、かなり執念深く追いかけています。しかも題材は、たった1行のベクター加算。a[i] + b[i] を並列に足すだけの、いちばん地味なCUDAプログラムです。地味だからこそ面白い。GPUという巨大な装置が、どれだけの段取りを踏んでこの1行を実行しているのかが、くっきり見えてきます。
nvcc は1つのコンパイラではなく、複数のコンパイラをまとめて動かす“司令塔”ですvadd<<<4096, 256>>>(...) のような記法は、ホスト側で引数を詰める「stub」に置き換えられますlibcuda.so、ドライバ、PCIe越しのGPUが連携して動きます元記事の出発点は、シンプルなCUDAコードです。3つの配列を用意して、GPUで足し算するだけ。
でも nvcc でコンパイルして実行すると、返ってくる c[0]=2.000000 の裏には、CPU命令が何千万回も走り、デバイスファイルが触られ、ioctl が900回も発行される、といった話が出てきます。ここがまず気持ちいい。GPU計算って、見た目は短いコードなのに、現実にはものすごく長い道のりを通っているんです。
私が面白いと思ったのは、この記事がGPUを“黒い箱”として扱わないところです。普通は「CUDAって速い」で終わります。でも著者は、その速さの正体を、コンパイラ、バイナリ形式、ドライバ、命令列、レジスタ配分まで降りて見にいく。こういう解剖は、知識欲がある人にはたまらないです。
nvcc は1本のコンパイラじゃないまず、CUDAプログラムを実機で動く形にするにはコンパイルが必要です。
ただし nvcc は、ありがちな「ひとつの巨大なコンパイラ」ではありません。実際には、いくつものコンパイラやツールを束ねて動かすドライバのような存在です。
記事では --keep を付けると、その途中生成物が見えると説明しています。そこには、たとえばこんなファイルが出ます。
sm_89 向けの cubinfatbinこの分解がとても大事です。
CUDAは「GPU用コードをそのまま1個のバイナリにする」わけではなく、ホストCPU向けのコードと、GPU向けのコードを別々に扱い、最後にまとめています。ここを知らないと、nvcc をただの不思議な黒箱だと思ってしまう。
PTXは、GPU実機そのものの命令ではありません。
著者はこれを virtual ISA と呼んでいて、要するに「実機依存をできるだけ薄めた中間表現」です。ここが実にCUDAらしいところです。
PTXの命令を見ると、たとえば配列アクセスのアドレス計算に何行も使っています。
i を計算し、境界チェックし、ポインタをグローバルメモリ用に変換し、sizeof(float)=4 を掛けて、足し合わせる。人間がCコードで一行に書いたことが、機械語寄りの世界ではいくつかのステップに割れて見えるわけです。
これはちょっとした「翻訳の現場」です。
人間語の「c[i] = a[i] + b[i]」を、そのままではGPUは理解できないので、コンパイラが超ていねいに言い換えている。こういう地味な変換を追うと、コンパイラという存在が少し身近になります。
PTXの次に待っているのが ptxas です。
これが、PTXを特定のGPUアーキテクチャ向けの SASS に変換します。記事では RTX 4090、つまり Ada 世代の sm_89 向けの SASS が例に出ています。
ここで見えるのは、もう“ふわっとした中間表現”ではなく、かなり実機っぽい命令列です。
S2R で special register を普通のレジスタにコピーし、IMAD で掛け算と加算をまとめ、LDG でグローバルメモリから読み、FADD で足し、STG で書き戻す。
見た目の印象としては、PTXよりずっと圧縮されている感じです。
著者が指摘しているように、PTXでは3命令必要だったアドレス計算が、SASSでは1命令に畳み込まれたりします。ここはコンパイラ最適化の旨味が一番わかりやすい部分です。私はこういう変形を見ると、コンパイラってただの変換機ではなく、かなり賢い“折りたたみ職人”なんだなと思います。
blockIdx.x や threadIdx.x は、特別なレジスタから来るSASSに出てくる S2R は “special register to register” の略です。
blockIdx.x や threadIdx.x のような値は、GPUがスレッドごとに持っている特別な情報なので、まずそれを一般レジスタにコピーしてから計算に使います。
これ、地味ですが重要です。
CUDAでは「どのスレッドが自分なのか」が計算の中心なので、そのための番号をGPUがちゃんと持っている。しかもそれを、人間の見るCコードの変数ではなく、ハードウェアに近い特別な仕組みで受け渡している。並列計算の現場感が出ます。
ここがかなり面白いところです。
カーネルの引数 a, b, c, n は、SASSの中で c[0x0][...] みたいな形で読まれます。これは constant bank 0 と呼ばれる領域です。
なぜわざわざ constant memory を使うのか。
理由は単純で、全スレッドに同じ値を配るのに向いているからです。配列ポインタや n は、グリッド内の全スレッドで共通です。ならば「全員に同じものを一回で配る」仕組みが効率的、というわけです。
元記事では、引数のオフセットが 0x160, 0x168, 0x170, 0x178 と固定で並んでいる点も説明されています。
つまり、ホスト側の launch stub が引数を詰める場所と、GPU側がそれを読む場所は、きれいに対応している。こういう対応表が見えると、GPU実行がぐっと“配線”っぽくなります。
<<< >>> が普通の関数呼び出しではなくなるCUDAの vadd<<<4096, 256>>>(...) という書き方、見慣れると不思議な記法ですが、あれは単なる文法の飾りではありません。
コンパイラはこれを見て、ホスト側に launch stub を自動生成します。
この stub は何をするかというと、引数をメモリ上のバッファに順番に詰めて、最後に CUDA runtime の launch API を呼びます。
要するに、見た目は関数呼び出しでも、実際は「GPUに渡すパケットを作っている」わけです。
私はここがかなり好きです。
人間向けには vadd(...) に見せておいて、裏では引数を丁寧に梱包し、別の世界へ配送する。昔の郵便みたいで、ちょっとロマンがあります。
GPUコードがバイナリに入っているだけでは、まだ何も起きません。
実行すると、ホスト側の CUDA runtime が libcuda.so を動的に開き、ドライバの世界へ入っていきます。
記事では strace を使って、その最初の openat まで確認しています。こういう観察の仕方が実に生々しい。
ふだんは「CUDAが動いた」で済ませるところを、実際には libcuda.so.1 が読み込まれ、context が作られ、デバイスにアクセスするための基盤が準備される、と見ていくわけです。
GPUはCPUの隣にいるようで、実際にはPCIeバスの向こう側にいます。
だからホストとGPUの間には、ただの関数呼び出しでは済まない“橋”が必要になる。この記事は、その橋のたもとから橋の先までを見せてくれます。
もうひとつ面白いのは、最終的な vadd バイナリの中に、いろいろなものが何重にも入っていることです。
cubinfatbinつまり、1つの実行ファイルの中に、複数の世界の成果物が同居しています。
GPUのプログラムは、思っている以上に“多層構造”です。単純なアプリのように見えて、実はかなり凝った配送システムになっている。ここに気づくと、CUDAの見え方が変わると思います。
個人的には、こういう記事の価値は「CUDAの細かい内部仕様がわかる」ことだけではないと思っています。
むしろ、「コンピュータは、見た目よりずっと段取りの塊だ」と実感できることが大きいです。
GPUの1回のカーネル起動は、派手な高速計算に見えます。
でも実際には、コンパイル段階で複数の表現に変換され、実行前にホスト側のstubが引数を詰め、ランタイムがドライバを呼び、ドライバがデバイスに命令を送り、GPUがようやくスレッド単位で計算を始める。しかもその間、命令やメモリ配置はかなり規律正しく整えられている。
こういう“下の層”を知ると、CUDAは単なる高速化テクニックではなくなります。
どこで抽象化が起き、どこで現実のハードウェアに合わせて折り畳まれるのか。その境目を見る話になる。そこがこの元記事のいちばんおいしいところだと思います。