CUDAカーネルを実行すると内部で何が起こるのか

 (fergusfinn.com)
2 ポイント 投稿者 GN⁺ 3 시간 전 | 1件のコメント | WhatsAppで共有
  • 単純なベクトル加算 CUDA プログラムでも、結果 2.000000 を得るまでに コンパイルパイプライン、ドライバ呼び出し、GPU コマンドキュー、ワープスケジューリング、メモリ階層、完了セマフォを経由する
  • nvcc はホストコードとデバイスコードを分け、cicc で PTX、ptxas で SASS を生成し、cubin と PTX を fatbin にまとめて Linux 実行ファイル内に格納する
  • vadd<<<4096, 256>>> の launch 構文はホスト launch stub に変換され、引数 da, db, dc, n は CUDA ランタイムと libcuda.so.1 を経由してドライバに渡される
  • GPU 実行は QMD、pushbuffer、GPFIFO、GP_PUT、doorbell MMIO 書き込みで開始され、RTX 4090 の 128 個の SM が 4096 個のブロックと 256 個のスレッド構成をワープ単位で実行する
  • このカーネルは float 加算 1 回あたり 12 バイトの転送が必要な 低い演算集約度 のため、Nsight Compute では 10.78μs、DRAM ピークの 79.65%、warp issue 5.17% となり、メモリ帯域幅に支配される

例のカーネルと観察範囲

  • 例のプログラムは vadd CUDA カーネルで 2 つの float 配列を加算し、3 つ目の配列に保存する
    • n = 1 << 20 として 1,048,576 個の float を処理する
    • launch 構成は vadd<<<4096, 256>>>(da, db, dc, n) で、4096 * 256 = n 個のスレッドを使う
  • RTX 4090 向けに nvcc -arch=sm_89 でコンパイルして実行すると、c[0]=2.000000 c[n-1]=2.000000 が出力される
  • この 1 行の結果にも、CPU 命令数千万回、device file、約 900 個の ioctl、メモリマップされた doorbell レジスタが関わっている

nvcc が実行ファイルを作る過程

  • nvcc --keep を使うと、コンパイルパイプラインの生成物を直接確認できる
    • vadd.ptx: cicc が生成したデバイスコードの PTX
    • vadd.sm_89.cubin: ptxas が生成したデバイスコードの SASS
    • vadd.fatbin: cubin と PTX をまとめた fatbin
    • vadd.cudafe1.stub.c: ホスト launch stub とカーネル登録コード
    • vadd.o: fatbin を含む最終的なホストオブジェクト
  • ホストコードはホストコンパイラで処理され、デバイスカーネル vadd は cicc と ptxas の段階を通る
  • PTX は仮想 ISA であり、型付きの無限個の仮想レジスタを使うため、実際のハードウェアレジスタ数を直接反映しない
    • 例の PTX には blockIdx.x * blockDim.x + threadIdx.x の計算、境界チェック、global load、float add、global store が含まれる
    • CUDA ポインタは基本的に generic pointer なので、cvta.to.global で global address に変換してから ld.global を使う
    • mul.wide.s32 は index を sizeof(float) である 4 バイト単位のオフセットに変換し、32 ビットから 64 ビットへ拡張する
  • SASS はアーキテクチャ固有の実命令で、RTX 4090 向け出力では PTX よりもさらに凝縮された形で現れる
    • S2R は SR_CTAID.X, SR_TID.X のような特殊レジスタを通常のレジスタへコピーする
    • PTX の mul.wide と add の組み合わせは、SASS では IMAD.WIDE に統合される
    • cvta 変換はアドレス指定の過程に吸収される
  • c[0x0][...] オペランドは、driver-managed な constant bank 0 を指す
    • ポインタ a, b, c は 0x160, 0x168, 0x170 に配置される
    • n は 0x178 に配置される
    • blockDim.x のような launch geometry と ABI 値も同じ bank にある
  • cubin は Linux 実行ファイルと同じコンテナ形式である ELF ファイルである
    • fatbinary は cubin と PTX を一緒にまとめる
    • この RTX 4090 では SASS が実際に実行されるが、PTX は他アーキテクチャでドライバが JIT コンパイルできるフォールバックとして含まれる
    • PTX は冗長なプレーンテキストなので、nvcc はデフォルトで圧縮する

ホストコードが launch を準備する方法

  • コンパイラフロントエンド cudafe++ は main より前に実行される隠れた constructor を挿入する
    • この constructor は埋め込まれた fatbinary を CUDA ランタイムに登録する
    • ホスト側の関数ポインタ vadd と fatbin 内のマングルされたデバイスカーネル名を結び付ける
  • vadd<<<4096, 256>>>(da, db, dc, n) 構文は、生成された host launch stub に変換される
    • da, db, dc, n はホストメモリ上の argument buffer に、それぞれオフセット 0, 8, 16, 24 で整列して格納される
    • これらのオフセットは、SASS が constant bank 0 から読む 0x160, 0x168, 0x170, 0x178 の位置に対応する
  • stub は __cudaLaunch を呼び出す際に、ホスト側のダミー vadd 関数アドレスを渡す
    • このアドレスは CPU 上で実行する関数アドレスではなく、ランタイム登録テーブルを参照するためのキーとして使われる
    • ランタイムは対応するデバイスシンボル名を探した後、クローズドソースのユーザーモードドライバ libcuda.so.1 に処理を渡す
  • 最初の GPU 呼び出し時、CUDA ランタイムは libcuda.so.1 を動的に開き、context を生成する
    • strace では /lib/x86_64-linux-gnu/libcuda.so.1 が開かれるのを確認できる
    • context には CPU が GPU と通信するための channel が含まれる
  • CUDA 12.2 以降では module loading はデフォルトで lazy である
    • 特定のカーネルが最初に launch されるまで SASS cubin のアップロードを遅延させる
    • CUDA_MODULE_LOADING で制御できる

GPU に作業を渡すコマンドキュー

  • GPU は CPU のように関数呼び出しを受けて entry point へ jump するわけではない
    • PCIe bus の向こう側から、ホストメモリ内にあるドライバの command stream を読む
    • cuLaunchKernel は完成した launch command をこの stream に入れ、GPU に通知する
  • 初回実行では、ドライバがカーネル SASS を GPU メモリへコピーする
    • code buffer を割り当てて SASS をコピーする
  • channel にはホスト RAM 上の 2 つの重要な構造がある
    • pushbuffer: ドライバが GPU command である method を書き込むメモリ領域
    • GPFIFO: pushbuffer span を指すポインタリングバッファ
  • GPFIFO entry は pushbuffer span の (base, length) を表す 2 個の 32 ビット word で構成される
  • GPU とドライバは 2 つのカーソルで作業の消費位置と生成位置を追跡する
    • GP_GET: GPU がどこまで消費したかを示す
    • GP_PUT: ドライバがどこまで生成したかを示す
    • どちらも USERD という per-channel 構造内にある
  • カーネル launch 時、ドライバは pushbuffer span に method を書き、GPFIFO entry がそれを指すようにした後、GP_PUT を進める
  • 現代の GPU では host engine がカーソルを常時監視していないため、doorbell が必要になる
    • GPU は process に小さな register window を mapping する
    • ドライバは channel の work-submit token を doorbell register に書き込む
    • host engine は doorbell を受け取ると GP_PUT を読み、GPFIFO entry と pushbuffer span を DMA で取り込む

QMD に含まれる実行情報

  • launch は SET_INLINE_QMD_ADDRESS_A/B と LOAD_INLINE_QMD_DATA の method burst で始まる
  • QMD(Queue Meta Data) は compute grid の launch descriptor である
    • grid と block サイズである 4096, 256 を含む
    • スレッドあたりのレジスタ数と shared memory 要求量を含む
    • プログラム開始アドレスと、カーネル引数を含む constant bank のアドレスを含む
    • 完了通知先も含む
  • ホスト stub がパックした引数はドライバによって constant bank にコピーされ、QMD にはその bank アドレスが記録される
  • QMD は GPU に対し、SASS の位置、並列プログラムの構成方法、完了シグナルの位置を知らせる
  • cuLaunchKernel は doorbell が鳴った時点で戻る
    • 呼び出しは非同期なので、CPU は GPU 作業の進行中も実行を続けられる

SM、ワープ、占有率

  • host engine は QMD を compute work distributor に渡す
    • この構成要素は GPU 全体に 1 つある
    • 線形な SASS instruction stream を各 SM に分配し、並列プログラムとして実行させる
  • 対象 GPU の GeForce RTX 4090 は 128 SM を使用する
    • launch は 4096 個の block と block あたり 256 スレッドで構成される
  • 各 SM はローカル instruction cache を持ち、active warp は program counter を維持する
    • Volta 以降ではスレッドごとの program counter と call stack を持つ Independent Thread Scheduling モデルがある
    • issue は依然としてワープ単位で行われる
  • 例のカーネルでは resource limit が block residency を決める
    • block あたり 256 threads = 8 warps
    • ptxas はスレッドあたり 16 個の register を確保する
    • register 基準では SM あたり 16 block が可能である
    • スレッド容量は SM あたり 1,536 active threads なので、1536 / 256 = 6 で 6 block しか置けない
    • したがって SM あたり最大 6 block、つまり 48 warp が resident 状態となる
  • SM は 4 個の processing block、すなわち sub-partition に分かれる
    • 48 個の resident warp は 4 個の sub-partition に均等分配される
    • 各 warp scheduler はフル状態で 12 個の active warp を管理する
    • 毎サイクル、eligible な warp を 1 つ選び、32 個の lane に次の命令を dispatch する

ワープが eligible 状態になる条件

  • GPU は CPU の out-of-order 実行のように、単一スレッドから大きな動的依存性を抽出しない
    • 多数の resident warp を持ち、stall が発生したら別の warp に切り替えてレイテンシを隠す
    • コンパイラが予測可能なタイミングを schedule し、hardware scoreboard が予測しにくい部分を処理する
  • 128 ビットの SASS instruction には、ptxas が書き込んだ control-code payload が入っている
    • 固定レイテンシ命令には static stall count が含まれる
    • yield hint は scheduler priority を譲るかどうかを示す
    • 可変レイテンシの操作にはワープごとの physical scoreboard barrier 6 個が使われる
  • 例の SASS 区間では、2 つの LDG.E が同じ scoreboard barrier B2 を set する
    • FADD は B2 を wait-on に持つ
    • 2 つの load が戻って barrier が clear されるまで、その warp は ineligible 状態になる
    • その間 scheduler は同じ sub-partition の別の warp を選ぶ
  • FADD から STG.E へ進む区間は固定レイテンシで処理される
    • FADD は stall=5 を持ち、R9 の結果が準備されるまで warp を数サイクル park する
    • 別の barrier は不要である
  • この control payload は nvdisasm のデフォルト出力では隠される
    • cuobjdump -sass の生の 128-bit encoding では、2 番目の 64 ビット word に含まれている
    • レイアウトは文書化されたものではなく、microbenchmarking によって再構成されたものだ

メモリアクセスと性能測定

  • warp が LDG.E を実行すると、32 スレッドがそれぞれアドレスを計算する
    • 例では連続した float 配列にアクセスするため、warp 全体で 32 * 4 = 128 bytes の連続ブロックを要求する
  • SM の load/store unit は request coalescing を行う
    • 32 個の 4 バイト要求を、4 個の 32 バイト sector request にまとめる
    • 連続アクセスでなければ、必要以上のデータを読む可能性がある
  • coalesced request はまず SM ローカルの L1 Data Cache を確認する
    • miss すると、crossbar interconnect を経由して 72MB の L2 Cache slice に向かう
    • L2 でも miss すると、memory controller と memory bus を通って GDDR6X VRAM に到達する
  • STG.E store も原則として逆方向の同じ経路をたどる
  • Nsight Compute の測定値は、このカーネルが memory-bound であることを示す
    • launch__grid_size: 4,096
    • launch__block_size: 256
    • launch__registers_per_thread: 16
    • launch__waves_per_multiprocessor: 5.33
    • sm__warps_active.avg.pct_of_peak: 82.77%
    • smsp__issue_active.avg.pct_of_peak: 5.17%
    • dram__throughput.avg.pct_of_peak: 79.65%
    • gpu__time_duration.sum: 10.78μs
  • カーネルは 演算集約度 が非常に低い
    • 2 回の 4 バイト load と 1 回の 4 バイト store、合計 12 バイト転送ごとに float add を 1 回実行する
    • DRAM read の観点では、8.4MB を 10.78μs で読み、およそ 780GB/s で、ピークの約 4/5 に達する
    • 4MB の出力 c は 72MB の L2 に収まるため、device-to-host copy が読み出すまで DRAM に flush されない

結果が CPU に戻るまで

  • kernel launch は doorbell を鳴らした時点で CPU に戻るため、GPU は完了を別途通知する必要がある
  • 4096 個の block がすべて retire すると、GPU は QMD に含まれる completion semaphore を post する
    • QMD の fence field は words 23–24 にある
  • default stream では cudaMemcpy(c, dc, ...) が kernel の後ろに置かれる
    • GPU copy engine は semaphore が上がるまで gated 状態になる
    • c はまだ 72MB の L2 に dirty 状態であるため、copy engine の read は DRAM 往復なしで L2 から処理される
    • データは PCIe を越えて host memory に移動する
  • copy が終わると、copy engine は自身の semaphore を post する
    • ホスト側の cudaMemcpy の待機が終了する
    • c は再び通常の host memory になる
    • printf は c[0] と c[n-1] を RAM から読み、stdout に出力する

launch の内部をのぞく方法

  • open kernel modules を読むだけでは、libcuda が closed-source なので一部の動作を直接確認しにくい
  • method write は syscall を経由せず、すでに mapping 済みの write-combined buffer に直接書かれるため、pushbuffer を見るにはメモリを読む必要がある
  • LD_PRELOAD shim で mmap をラップし、/dev/nvidia* で mapping された領域を記録できる
    • test program が launch 直後に shim の dump 関数を呼び出せば、mapping された pushbuffer を出力できる
    • dump は SET_INLINE_QMD_ADDRESS_A に対応する method burst を探す
  • pushbuffer method header は opcode、payload count、subchannel index、register offset を bit field で持つ
    • 0x0318 は SET_INLINE_QMD_ADDRESS_A
    • 0x0320 + i * 4 は LOAD_INLINE_QMD_DATA(i)
    • dump では count 66 の increasing-method burst が見え、2 個の address word と 64 個の QMD word、合計 256 バイトの QMD が inline で載る
    • QMD 内の word 12 は 0x1000、word 18 は 0x100 で、launch の 4096 と 256 に対応する
  • ドライバのセットアップは ioctl で行われる
    • one-kernel program では strace が 948 個の ioctl を記録する
    • その大半は one-time setup である
    • 主要な file descriptor は /dev/nvidiactl と /dev/nvidia-uvm である
    • NVIDIA resource manager ioctl magic byte は 0x46、すなわち 'F' である
    • command number 0x2A は NV_ESC_RM_CONTROL、0x2B は NV_ESC_RM_ALLOC と解釈される
  • nvcc --keep で生成される vadd.cudafe1.stub.c では、startup registration code も確認できる
    • __attribute__((__constructor__)) が付いた関数が main より前に実行される
    • __cudaRegisterBinary と __cudaRegisterEntry によって、host function pointer vadd と device entry point _Z4vaddPKfS0_Pfi が結び付けられる

関連記事

1件のコメント

 
GN⁺ 3 시간 전
Hacker News のコメント

  • 興味深い記事で、デフォルトストリームのセマフォの説明も面白かった
    CUDA がコマンド同期を暗黙に処理してくれて、並列コマンドはストリームを通じて選択的に使えるようにしている点が良い
    最初から同期の複雑さをすべてユーザーに押し付ける Vulkan とは対照的

  • ハードウェア側については一部公開ドキュメントがある
    メソッドのドキュメントや QMD 形式を探すために、必ずしもカーネルソースを読む必要はない
    https://github.com/NVIDIA/open-gpu-doc/blob/master/classes/c... を参照

  • とても有用だった
    特に doorbell と QMD の部分が、CUDA の実行構文が実際に GPU へ投入される内容とどうつながるのかを示していて、いちばん役に立った
    たいていの説明はカーネル、ブロック、ワープあたりで止まるが、この記事は CPU→ドライバ→GPU の経路をずっと追いやすくしてくれる

  • 制御コードは記事で説明されているより少し複雑
    実際には制御ワード内のビットというより、テーブル参照に近い

  • 今では、カーネルを最適化してより高速に動かすことを主業務にする会社がある
    そうした会社がいつか、それを非常にうまくこなすオープンソースライブラリに押されるのか気になる
    Nvidia ならいつでもそういうものを出せそうにも見える
    あるいは、大手プロバイダが推論速度を高めるための moat としてこれらの会社を買収し、むしろうまくいく可能性もある

    • 短期的には人材獲得型の買収がかなりあり得そう
      ただ、kernelbench のような関連ベンチマークでモデルが進歩しているのを見ると、より汎用化された解法も結局は出てこざるを得ないと思う
      問題は、新しいハードウェア世代ごとに、既存モデルが見たことのない制約や機能が頻繁に出てくること
      たとえば Blackwell の tcgen05 も、かつては分布外の例だった
      モデルがよりうまく汎化し始めれば致命的な障壁ではないかもしれないが、少なくとも今はまだ障害になっている
      [1] https://kernelbench.com/
    • CUDA を大規模に動かすと、Nvidia ドライバとライブラリのバグへの対処に、気分が悪くなるほど多くのエンジニア時間が取られる
      Nvidia のライブラリにさらに依存することを期待している人はあまり見たことがない
    • おそらくそうはならないと思う
      ワークロードの細部、つまり正確なパラメータ、メモリ内でのデータ表現、値の範囲などによって、最適化戦略が大きく分かれるから

  • HPC の修士を終えたばかりで、CUDA、MPI+CUDA、OpenCL の授業を受けたが、授業の前にこういう記事を読んでいたらずっと役に立ったと思う
    特に、ワープが実行可能であるという意味を扱った部分の前後が良かった

  • まず、いろいろな細部をよく掘り下げた良い記事
    ただし CUDA の runtime API を経由しなければ、ユーザー空間にある多くのブードゥーめいた部分は消える
    ドライバ API を使い、カーネルソースを文字列として受け取って NVIDIA のランタイムコンパイラでコンパイルすれば、何が起きているのかをよりよく見られる
    すべてではないが、かなりの部分が透明になる
    より「原始的な」バージョンはここにある:
    https://github.com/NVIDIA/cuda-samples/tree/master/cpp/0_Int...
    同じ内容をずっと読みやすく、それでも完全に透明な現代的 C++ API の形で見るなら、これを見るとよい:
    https://github.com/eyalroz/cuda-api-wrappers/blob/master/exa...
    私の CUDA API wrappers ヘッダーオンリーライブラリのサンプルプログラム

    • ドライバ API は CUDA カーネルをホットリロード可能なシェーダーのように扱えるので良い
      実行中にコードを変えながら開発できるので楽しい

  • ベアメタルで？