libclang/libLLVM で CUDA 言語相当の JIT コンパイル機能を自前 C++ アプリに組み込んで CUDA(PTX) 実行するメモ

背景

• clang/LLVM を使って, C/C++/CUDA JIT アプリを自前 C++ アプリに組み込みたい
• CUDA だと, NVRTC などはあるがシステムに nvcc などがインストールされている必要があり, アプリ単体配布で完結するというのができない
• prebuilt package だと, RTTI/例外などのコンパイルオプションが違ったりする + STL が違ったりでアプリへのリンク + JIT コンパイルがうまくいかないため, clang/LLとVM をソースからビルドする必要がある

目的

nvcc など SDK 不要で, nvidia ドライバだけインストールされている環境で CUDA(C/C++ like) カーネルを JIT コンパイルしてうごかす.

• OpenCL 1.2 のように, device 単体でのカーネル記述を考えます
  • host と device が混在するような記述は考えない
• libc 関連や C++ STL のサポートは今回は考えないことにします
  • 使えるのは assert, printf, math 関数くらい
• とりあえずは Linux のみ. libc, C++ STL を考えればそれほど難しくなく Windows 環境(llvm-mingw, MSVC(clang-cl))も対応できると思われます.

nvcc 自体は最近は clang/llvm ベースなので(昔は Open64 compiler ベースだった), clang/llvm の repo にあるので大体 nvcc がサポートする CUDA 言語の対応ができると思われます.

以前は clang/llvm に CUDA 相当を取り入れる gpucc という動きもありましたが, オフィシャル(NVIDIA CUDA SDK)に最終的にはマージされたのかしらん?

成果物

途中ですが,

にあります

llvm/clang コンパイル時間のメモ

llvm/clang, 通常の単体ビルド(MinSizeRel target)だと, Ryzen9 3950X + NVMe SSD(PCI Gen4)でおよそ 7 分ほどです.

clang/llvm 組み込み方法

いろいろ一式ある llvm-project repo を使います.
(flang など使わないものもあるが, 個別に repo 管理めんどいので一式を引いたほうが効率がよい)

submodule で追加

の Cmake を参照ください. clspv の cmake を参考にしています.

最近は clang/llvm も cmake サポートがこなれてきたようで,
cmake で add_subdirectory で追加するだけです. あとは使うライブラリ名(e.g. libLLVM***.a など)をうまく列挙すれば組み込めます.

include の設定は LLVM/clang 側 cmake での設定に含まれていないので, 明示的に指定します.

ただし, この方法だと, 自前 C++ コードを変えてリビルドすると llvm/clang の各ビルドターゲットのチェックに数秒かかる + llvm/clang の C++ API 自体がコンパイルが遅い + リンクに時間がかかるで, 開発するときにはかなり不向きです
(CI でのビルドや, 最終リリースビルド用に使うのがよい)

• dll(shared lib)版を使う(リンクを早くするため. ただ lld 使うなら .a リンクでもあまり変わらないかも)
• python などのバインディングでひとまず記述して, API の挙動など確認してから C/C++ で書き直し

がよいでしょう.

C API を使いところですが, clang のほうはソースコードをパースしたりくらいまでの機能しか expose されおらず, libclang の C API では JIT 機能を実現できないです(ELF バイナリやコンパイル後の asm 取得など)はありません.
(LLVM C は行けるかもなので, clang を外部呼び出しで一度 bc にするという手もあるかもだが...)

libclang

ややこしいですが, clang の cmake では, libclang がライブラリのターゲットとして定義されていて clang C API なども含んでいるライブラリ一式になります
(C API + 各 clangAST, clangBasic, ... などの個別 clang C++ ライブラリをまとめている)

したがって, clang を add_subdirectory で追加している場合は, libclang を指定します.

target_link_libraries(${MINIJIT_BUILD_TARGET} PRIVATE libclang)
# ライブラリ名は `clang` ではない!

ビルド済み(install 済み)の libclang.so とリンクする場合は通常どおり clang を使います.

libLLVM.so, libLLVM-C.so(macOS のみ)

.a だと LLVM はいろいろライブラリあってめんどいので, ひとまとめになっている libLLVM.so を作るのがよいです.

LLVM_BUILD_LLVM_DYLIB On
LLVM_BUILD_LLVM_C_DYLIB On(macOS のみ)

を指定します.

libLLVM.so, libclang.so 両方リンクではエラーが出る

ただし, libLLVM.so, libclang.so 両方をリンクすると, 実行時の dll ロード時に

CommandLine Error: Option 'help-list' registered more than once! 
LLVM ERROR: inconsistency in registered CommandLine options

というエラーが出てしまいます

何かしら同じシンボルか初期化ルーチンが両方の .so に追加されてしまっているようです.

LLVM 側だと .a のリストがいっぱいあって面倒なので, とりあえず clang 側を .a リンクで解決します.

clang/LLVM をビルドしてから自前アプリに追加

開発者にはこちらを推奨します.

minijit の scripts/build-llvm-clang-linux.sh を参照ください.

以下にあるように, デフォルトだと RTTI/EH off でビルドされるので, 自前アプリの C++ 構成やコンパイラと合わせてビルドするようにしましょう.

ビルドオプション

clang/LLVM のビルドで気をつけるのは以下です.

• RTTI, EH(Exception Handling)
  • LLVM ではデフォルトで off です. 有効にしたい場合は cmake ビルドオプションで有効にする必要があります
• terminfo(Unix 系のみ)
  • curses? で色付きでコンパイルメッセージ(エラーメッセージなど)を出す用でしょうか. デフォルトではシステムに libterminfo などある場合に有効になりますが, 配布先のシステムでインストールされていないとリンクエラーなり実行時 dll シンボル未解決エラーになるので, 基本 disable しておいたほうがよいでしょう.
• lld リンカ
  • lld linker がインストールされている場合(ld.lld コマンドがある), -DLLVM_USE_LINKER=lld を指定するとビルド(リンク)が早くなります(linux だとデフォルトは bfd). 出来上がるライブラリ .a(.so) に違いはありません.

その他ビルドオプションは https://llvm.org/docs/CMake.html#llvm-specific-variables を参照ください.

CUDA JIT(PTX target)の準備

llvm/Config/Targets.def に NVPTX があるか確認しましょう.

NVPTX があれば,

LLVMInitializeNVPTXTarget() などの関数が定義されます.

cuda 関連のヘッダ

clang/llvm をインストールすると,

<dist>/lib/clang/12.0.0/include/__clang_cuda_***.h

に CUDA 関連のヘッダがあります(builtin 定義など)

__device__ などは, 本来は CUDA SDK 側のヘッダを呼んで解決します(crt/host_defines.h あたり)が, CUDA SDK インストールを想定していない場合は自前定義をでっち上げます.

#define __device__ __attribute((device))

など.

<clang>/test/SemaCUDA/Inputs/cuda.h に定義のサンプルがあります.

CUDA コードを PTX に変換

-x cuda を指定します.

__syncthread() などは <clang>/include/clang/Basic/BuiltinsNVPTX.def に宣言があります.

References

• Using libclang to Parse C++ (aka libclang 101) https://shaharmike.com/cpp/libclang/
• Getting started with the (recent) LLVM C JIT API https://www.owenstephens.co.uk/blog/2018/09/25/getting-started-with-the-newer-llvm-c-api.html
• Halide の PTX CodeGen: https://github.com/halide/Halide/blob/master/src/CodeGen_PTX_Dev.cpp

TODO

• 似たようなやりかたで OpenCL for C++ に対応する(一応すでに clspv などあるが...)