概要
- CUDAプログラムが ベクトル加算 を実行する流れを解説
- コードの コンパイル過程 から GPU実行 までを詳細に追跡
- PTX と SASS への変換、バイナリ構造の説明
- ホストからGPUへの カーネル起動手順 を解説
- ドライバとハードウェアの連携 の仕組みも紹介
CUDAベクトル加算プログラムの全体像
- CUDAプログラム例： 2つのfloat配列の要素ごと加算
- カーネル
__global__ void vadd(...)内で スレッドごとに加算結果を格納 main()で 配列メモリ確保・初期化・デバイス転送・カーネル起動・結果取得 を実施vadd<<<4096, 256>>>で 1,048,576スレッド を生成、各要素に1スレッド割当- RTX 4090 などのGPUで実行可能な構成
nvccによるコンパイルと生成物
- nvcc は複数のコンパイラを連携させるドライバプログラム
--keepオプションで 中間生成物 をディスクに保存可能vadd.ptx：PTX（仮想ISA、cicc生成）vadd.sm_89.cubin：SASS（実機コード、ptxas生成）vadd.fatbin：PTXとcubinを束ねたファットバイナリvadd.cudafe1.stub.c：ホスト側起動スタブ・カーネル登録vadd.o：最終ホストオブジェクト（fatbin埋込済み）
PTX（仮想ISA）の特徴
- 無限個の型付きレジスタ を仮想的に利用可能
- アドレス計算は複数命令に分解される
- CUDAポインタはデフォルトで「generic」型、cvta命令でグローバルメモリに変換
- PTXは デバイス非依存 の記述
SASS（実機コード）の特徴
- PTXから SASS（実デバイス命令列） へ変換
- 実レジスタ数 に制約される
- アドレス計算や特殊レジスタ（blockIdx.x, threadIdx.x）の扱いがハードウェアに最適化
- カーネル引数や起動ジオメトリは constant bank 0 に格納
- 例：a, b, c, nの各ポインタや値は固定オフセットに配置
fatbin・バイナリ構造
- fatbin はcubin（SASS）とPTX（テキスト）を束ねて格納
- PTXは将来的な互換性のために同梱
- 未対応アーキテクチャではPTXからJITでSASS生成
- fatbinは ELF形式 でホストバイナリに埋め込まれる
.nv_fatbinセクション等に配置
ホストからGPUへのカーネル起動の流れ
-
cudafe++ が隠れたコンストラクタを挿入し、fatbinをCUDAランタイムに登録
-
カーネル呼び出し
vadd<<<4096, 256>>>は 自動生成されたスタブ関数 に置き換え -
スタブ関数は 引数をバッファにパック し、
__cudaLaunchを呼び出し- da, db, dc, nの各引数は16バイトアライメントで配置
- この配置はSASSがconstant bank 0から読むオフセットと一致
-
__cudaLaunchは ホスト側vaddのアドレス をキーに、デバイス側カーネル名を対応付け -
CUDAランタイムは libcuda.so.1 （ユーザーモードドライバ）を動的にオープン
- 初回呼び出し時に GPUコンテキスト を作成
- コンテキストにはデバイス通信チャネル等の情報が含まれる
GPUへのカーネル展開と実行
- CUDA 12.2以降、 モジュールロードは遅延（LAZY） がデフォルト
- 初回カーネル起動時にのみSASSをGPUメモリへコピー
- libcudaの下層には カーネルモードドライバ（nvidia.ko）
- ホストとGPU間は ioctl とデバイスファイルで通信
- cuLaunchKernelは コマンドストリーム をPCIe経由でGPUに送信
- GPUは関数呼び出しやスタックではなく、 命令ストリームを逐次実行
この流れにより、 シンプルなCUDAプログラム でも、コンパイル・バイナリ構成・カーネル起動・デバイス実行まで、多層的な仕組みが連携し、最終的に GPU上で高速な並列演算 が実現される。