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)に最終的にはマージされたのかしらん?

http://wujingyue.github.io/docs/gpucc-talk.pdf

## 成果物

途中ですが,

https://github.com/syoyo/minijit

にあります

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

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

## clang/llvm 組み込み方法

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

https://github.com/llvm/llvm-project

### submodule で追加

https://github.com/syoyo/minijit

の 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
```

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

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

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

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

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

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

https://github.com/syoyo/minijit/blob/master/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 などあるが...)