OpenMP Offloading for NVIDIA GPU

3 min read読了の目安(約3400字

OpenMP Offloading を使うためではなく、どうやって動作しているのかを理解するためのメモ

Programming model by NVIDIA

  • 最適化されたライブラリの使用 (cuBLAS, cuDNN等)
    • 特別なコンパイラ無しで使用できる、用意されたコードしか GPU 上で動作させられないので柔軟性は低い
  • OpenACC 言語拡張を用いて既存のコードを GPU 向けに自動的に変換する
    • for 文に #pragma acc を追加すれば GPU 上で動作するようになる
  • 独自言語拡張 CUDA/C++ を用いる
    • NVIDIA が提供している全ての機能にアクセスできる
  • C++17 Parallel Algorithm, Fortran 2008 concurrent を GPU で実行する
    • nvc++, nvfortran (NVIDIA HPC SDK) で標準の C++ / Fortran をコンパイルする

How NVIDIA GPU / CUDA works?

  • 何かしらの方法で GPU 上で実行されるアセンブラ (PTX) を生成し、それを GPU 上に転送し、起動する必要がある
  • ハードウェアとしての GPU はカーネル (Linux, Windows) にロードされた NVIDIA ドライバが面倒を見る
    • macOS はカーネルドライバが提供されないので全く使用できない
  • ユーザー空間からカーネルへの命令は CUDA Runtime を経由して行われる
    • CPU の場合の C Runtime (glibc, CRT) のようなもの
  • CUDA Runtime には高水準な Runtime API (libcudart.so) と低水準な Driver API (libcuda.so) がある
    • 基本的に Runtime API は Driver API を呼び出して動いているはず?
  • CUDA/C++ compiler (nvcc) はプログラム中の __device____global__ 修飾されている関数を PTX にコンパイルし、カーネル呼び出し構文 kernel<<<block, thread>>> を CUDA Runtime の呼び出しに置き換える
    • nvcc の機能をライブラリとして提供する NVRTC というものも存在する
  • PTX さえ生成出来れば特別なコンパイラ無しに、それを CUDA Runtime を用いて GPU 上で実行させることが出来る
    • PTX ISA reference を読みながら手で書く
    • GCC, LLVM は PTX をターゲットとしてコード生成する事が可能
      • ビルド時にNVPTX が有効になっている clang があれば clang --target=nvptx64-nvidia-cuda -S test.cpptest.s に PTX が出来る
      • 例えば RustJulia, あるいは Python から LLVM IR を経由して PTX へコンパイルすることが出来る
      • OpenMP Offloading もこの仕組みを使う

OpenMP Offloading

  • GPU等のアクセラレータを OpenMPコンパイラ から使えるようにしたもの
    • OpenMP とは C/C++/Fortran 向けスレッド並列化用言語拡張 (#pragma omp)
    • OpenMP Offloading は CUDA/C++ コンパイラ (nvcc) と同じようにユーザーコードから GPU 用のアセンブリを生成し、それを GPU に転送・実行をするための処理を記述するための言語拡張
  • OpenMP OffloadingOpenACC は別の言語拡張
    • コンパイラの裏側では同じ実装だったりする
  • OpenMP はコンパイラでの言語拡張なので各コンパイラ毎に実装する
  • GCC, LLVM 共に NVIDIA GPU の制御には CUDA Runtime を使う
    • コンパイラ (gcc, clang) はデバイスコードを PTX にコンパイルしてオブジェクトファイルに埋め込み、ランタイム (libgomp, libomp) が実行時にそれを CUDA Runtime で GPU 上に展開・実行させる
    • リンカはホストコードのリンクに加えて、デバイスコード (PTX) をリンクする必要がある
      • ここは標準化されている?
    • ld は PTX をリンク出来ないので ptxas (CUDA の一部) を使うか別実装を使う
  • GCC は PTX をオブジェクトに埋め込んで nvptx-tools を使ってリンクする
  • LLVM は調査中(挙動が結構違うので多分中身も結構違う)

この記事に贈られたバッジ