CUDAカーネルを実行するとどうなるか？

CUDAカーネルを実行するとどうなるか 2026年6月29日 ・ 35分読了 ・ カバー: サロモン・ド・コーズ作「レ・レゾン・デ・フォース・ムーヴァント」（1615年）より、ピンシリンダー式水オルガンの版画。 ここにシンプルなCUDAプログラムがあります。2つのベクトルを加算します。 ```c __global__ void vadd(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]; } int main() { int n = 1 << 20; // a million floats (1,048,576) size_t bytes = n * sizeof(float); float *a = (float*)malloc(bytes), *b = (float*)malloc(bytes), *c = (float*)malloc(bytes); for (int i = 0; i < n; i++) a[i] = b[i] = 1.0f; float *da, *db, *dc; cudaMalloc(&da, bytes); cudaMalloc(&db, bytes); cudaMalloc(&dc, bytes); cudaMemcpy(da, a, bytes, cudaMemcpyHostToDevice); cudaMemcpy(db, b, bytes, cudaMemcpyHostToDevice); vadd<<<4096, 256>>>(da, db, dc, n); // 4096 * 256 = n threads, one per float cudaMemcpy(c, dc, bytes, cudaMemcpyDeviceToHost); printf("c[0]=%f c[n-1]=%f\n", c[0], c[n-1]); } ``` RTX 4090用にコンパイルして起動すると、1+1=2が百万回正しく計算されます（すべては確認していません）。 ```bash $ nvcc -arch=sm_89 -o vadd vadd.cu && ./vadd c[0]=2.000000 c[n-1]=2.000000 ``` これを実行するには、数千万のCPU命令、いくつかのデバイスファイル、900のioctl、そして1つのメモリマップド・ドアベルレジスタが関与しています。この記事では、この1つのカーネルがコードからワープ、そして答えが返されるまでの過程を追っていきます。（余談ですが、この記事はエージェントによって生み出された「読みやすさの変遷」の一例です。好奇心と（機械による）持続力があれば、コンピュータについて知ることができないことはほとんどありません。AIが私たちが何を知るのに役立つかについての読みやすさの示唆についての興味深い議論はこちらです。） nvccでプログラムをコンパイルする デバイスが実際に読み取れるものにこのCUDAプログラムを変換する方法から始めるべきでしょう。そのためにはコンパイラが必要です。実際には、多くのコンパイラが必要です。nvccは、いくつかの他のコンパイラを実行し、その出力を結合するドライバープログラムです。--keepオプションを渡すと、パイプライン全体がディスク上に残され、読み取ることができます。 ```bash $ nvcc --keep -arch=sm_89 -o vadd vadd.cu && ls ... vadd.ptx # device code as PTX (from cicc) vadd.sm_89.cubin # device code as SASS (from ptxas) vadd.fatbin # cubin + PTX, bundled (from fatbinary) vadd.cudafe1.stub.c # host launch stub + kernel registration vadd.o # final host object, fatbin embedded ... ``` ホストコードはホストコンパイラに送られます。デバイスコード（vadd）はさらにいくつかのステップを要します。LLVMベースのコンパイラであるciccがそれをPTXに変換し、次にptxasがPTXをSASSに変換します。PTXは仮想ISAです。無限の型付きレジスタを持ち、ハードウェアが実際にいくつのレジスタを持っているかという概念はありません。以下はvaddの（省略された）PTXの本体です。 ```bash $ cat vadd.ptx ... mad.lo.s32 %r1, %r3, %r4, %r5; // set register r1 to ctaid*ntid + tid setp.ge.s32 %p1, %r1, %r2; // set predicate p1 if i >= n @%p1 bra $L__BB0_2; // if out of bounds, skip to exit cvta.to.global.u64 %rd4, %rd1; // convert generic pointer %rd1 to a global address, store in %rd4 mul.wide.s32 %rd5, %r1, 4; // multiply r1 by 4, store the result in %rd5 add.s64 %rd6, %rd4, %rd5; // add %rd4, %rd5, result in %rd6 ld.global.f32 %f2, [%rd6]; // load a[i] into %f2 ... add.f32 %f3, %f2, %f1; // add %f1 and %f2, result in %f3 st.global.f32 [%rd10], %f3; // store c[i] = ... in global memory ``` 仮想レジスタは%rd1～%rd10、%f1～%f3のように見えます（プレフィックスは型を示します：%rは32ビット整数、%rdは64ビット、%fは32ビット浮動小数点数、%pは1ビット述語です）。PTXは思ったよりも「回りくどい」です。例えば、%rd6に1つのアドレスを形成するのに3つのPTX命令が必要です。これは、PTXがデバイスに依存しないためです。なぜ3つなのでしょうか？CUDAポインタはデフォルトで「汎用」であり、グローバル、共有、またはローカルメモリを参照できることを意味します。cvta.to.globalはポインタがグローバルウィンドウに存在することをアサートするため、後でより安価なld.globalを使用できます。次に、mul.wide.s32はインデックスiに4（sizeof(float)）を掛けて32→64ビットに拡張し、バイトオフセットに変換します。add.s64はそれにベースポインタを加算します。 次に、ptxasはデバイスに依存しないPTXを、アーキテクチャに特化したSASSに変換します。生成されるSASSは異なります。 ```bash $ cuobjdump -sass vadd /*0000*/ MOV R1, c[0x0][0x28] ; // set up the stack pointer (ABI; unused here) /*0010*/ S2R R6, SR_CTAID.X ; // R6 = blockIdx.x /*0020*/ S2R R3, SR_TID.X ; // R3 = threadIdx.x /*0030*/ IMAD R6, R6, c[0x0][0x0], R3 ; // i = ctaid*ntid + tid /*0040*/ ISETP.GE.AND P0, PT, R6, c[0x0][0x178], PT ;// P0 = (i >= n) /*0050*/ @P0 EXIT ; // if so, exit /*0060*/ MOV R7, 0x4 ; // load literal 4 (sizeof(float)) into R7 as multiplier /*0070*/ ULDC.64 UR4, c[0x0][0x118] ; // uniform load of a driver-provided system value /*0080*/ IMAD.WIDE R4, R6, R7, c[0x0][0x168] ; // &b[i] /*0090*/ IMAD.WIDE R2, R6, R7, c[0x0][0x160] ; // &a[i] /*00a0*/ LDG.E R4, [R4.64] ; // b[i] /*00b0*/ LDG.E R3, [R2.64] ; // a[i] /*00c0*/ IMAD.WIDE R6, R6, R7, c[0x0][0x170] ; // &c[i] /*00d0*/ FADD R9, R4, R3 ; // a[i] + b[i] /*00e0*/ STG.E [R6.64], R9 ; // c[i] = ... /*00f0*/ EXIT ; ``` S2R行は何をしているのか S2Rは「特殊レジスタからレジスタへ」を意味します。これは、ハードウェアがスレッドごとに維持する特殊レジスタ（ここではSR_CTAID.X（ブロックのインデックス、blockIdx.x）とSR_TID.X（ブロック内のレーンのインデックス、threadIdx.x））を通常のレジスタにコピーし、IMADがそれに対して算術演算を実行できるようにします。10個ほどの仮想レジスタが7個の物理レジスタに集約されました（ncuの報告ではlaunch__registers_per_thread = 16です。逆アセンブルではR9までしか名前が付けられていませんが、アロケータはABIとアライメントのためにさらにいくつか予約しています）。2つのmul.wideとaddのシーケンスは、単一のIMAD.WIDEに統合されました。cvta変換はなくなり、アドレス指定に吸収されました。c[0x0][…]オペランドは定数バンク0にあり、ドライバーが管理する小さな領域にあります。これらはカーネルの引数（ポインタa、b、cとサイズn）、および起動ジオメトリです（blockDim.xは0x0にあります）。バンクを埋めるのは、起動時にドライバーがGPUに渡すQMDと呼ばれる構造体の仕事であり、カードに起動が届いたときにそれについて説明します。 なぜ引数が定数バンク0に置かれるのか、そしてどこに これらが定数メモリにあるのは、これがブロードキャスト読み出しだからです。グリッド内のすべてのスレッドが同一のポインタを必要とし、定数キャッシュは32レーンすべてに一度にサービスを提供できます。レイアウトは固定されています — 0x160、0x168、0x170はポインタa、b、cであり、0x178はnで、起動ジオメトリは0x0に並んでいます（blockDim.x）。バンク0には、MOV R1, c[0x0][0x28]がエントリでロードするスタックベースのようなABIパラメータも保持されています。ホストスタブが起動のために引数をパッキングするときにも、これらの同じオフセットを見ることになります。このSASSを保持する「cubin」ファイルはELFファイルです — Linuxが通常の実行ファイルや共有ライブラリに使用するのと同じオブジェクトファイルコンテナです（cuobjdump -elfは、シンボルテーブル、マシンコードを保持する.text.vaddセクション、および.nv.callgraphのようなCUDA固有のセクションを示します）。fatbinary実行ファイルは、cubinをPTXと一緒に単一の「fatbin」にバンドルし、結果をcuobjdumpで確認すると、バイナリに埋め込まれたfatbinには両方が含まれていることがわかります。 ```bash $ cuobjdump vadd ... Fatbin elf code: arch = sm_89 # the SASS we just read Fatbin ptx code: arch = sm_89 compressed # the PTX, shipped too ``` SASSはこの4090で実際に実行されるものですが、PTXは前方互換性のフォールバックとして同梱されます。このバイナリを、cubinがカバーしていないアーキテクチャのGPUに持っていくと、ドライバーはロード時にPTXを新しいSASSにJITコンパイルできます。最後に、そのfatbinはホスト実行可能ファイル内にネストされており、readelf -Sで独自のセクションを占めていることがわかります。 ```bash $ readelf -S vadd ... [18] .nv_fatbin PROGBITS ... [19] __nv_module_id PROGBITS ... [29] .nvFatBinSegment PROGBITS ... ... ``` nvccが吐き出すvaddバイナリは、ホストコード、Ada SASSを含む完全なELFオブジェクト、およびPTXのコピーを含む単一の実行可能ファイルです。PTXは冗長なプレーンテキストであるため、nvccはバイナリサイズを小さく保つためにデフォルトで圧縮します。ドライバーはそれを解凍するだけです。