NVRTCライブラリを使う
CUDA7.0RCの新機能の一つに、Runtime Compilationというのがある。
従来、cuModuleLoad()などでGPU用バイナリをロードして使う際には、nvccを実行してC/C++のソースからPTXなりのバイナリを生成する必要があった。CUDA 7.0RCのRuntime Compilationは、これをOpenCLライクな関数呼び出しでソース→バイナリへの変換を行うライブラリである。
予め静的なコードを書いてコンパイルしておく、というレベルの話であれば一向に構わないのだが、PG-Stromの場合はGPUのコードは実行時に動的生成されるので、GPUコードのビルド時間は割と性能影響の大きいファクター。あまりチンタラやっている訳にはいかない。
以下の図はCUDA Toolkitのドキュメントにあるnvccの処理フローを示した図だが、nvccに何を食わせるかによっても処理フローが変わってくる。つまり、ビルド時間に影響がある。
まず考えねばならないのは、入力となるソースコードは何であるか。
.cu形式と.gpu形式が選択肢ではあるが、どちらも一長一短である。
- .cu形式の場合、先ずホストコード/デバイスコードの分離を行い、さらにC/C++で二通りのビルドプロセスが走る。要はもっさりしている。その一方で、CUDA関数が利用できるよう適切なファイルを勝手にインクルードしてくれたりと、デプロイを考える上では非常に楽。
- .gpu形式の場合、予めデバイスコードとして分離されたものに対してビルドを行うので、ビルド時間は非常に短い。ただ、.cu形式であればnvccが勝手に解決してくれたようなデプロイにまつわるあれやこれやを自分で指定してやる必要がある。正直、これをNVIDIA謹製ツール以外が行うのは違うと思う。
という悩みがあった訳だが、CUDA 7.0RCでサポートされたNVRTCはこの辺の問題をきれいさっぱり解決してくれた。
まず最初に、cstring形式のソースコードを与えて nvrtcProgram オブジェクトというものを作成する。
nvrtcResult nvrtcCreateProgram(nvrtcProgram *prog,
const char *src,
const char *name,
int numHeaders,
const char **headers,
const char **includeNames)次に、このnvrtcProgramオブジェクトをビルドする。
nvccのオプションを全てサポートしている訳ではないが、いくつかのコンパイルオプションを指定する事もできる。
nvrtcResult nvrtcCompileProgram(nvrtcProgram prog,
int numOptions,
const char **options)で、ビルドが終わった後のnvrtcProgramオブジェクトからは、PTXイメージやビルドログを取り出す事ができる。
nvrtcResult nvrtcGetPTX(nvrtcProgram prog, char *ptx) nvrtcResult nvrtcGetProgramLog(nvrtcProgram prog, char *log)
当初、このAPIを見た時に、nvrtcCompileProgram()関数の呼び出しが同期呼び出しになってしまうので、ソースのビルド中にテーブルスキャンを先に進められないのはイカガナモノカ…と思ったのだが、よく考えたら、こんなものはDynamic Background Workerを使えば何とでもできる。
例えば、エグゼキュータの開始時に nvrtcCompileProgram を呼び出して（あるいは、従来はnvccコマンドを実行して）、GPUコードのビルドを行うDynamic Background Workerプロセスを立ち上げておく。
この時点ではまだ、GPUで命令を処理するためのバイナリは生成されていないので、PostgreSQL本体側はせっせとテーブルのスキャンに励み、データをDMAバッファに積んでおく。
で、しばらくするとBgWorker側でGPUコードのビルドが完了しバイナリが生成されるので、BgWorkerはSetLatch()を使ってバックエンド側にビルドが完了した事を通知する。
この時点で動作モードが切り替わり、今までメモリに読み込んできたデータブロックのDMA転送とGPUカーネルの起動が始まると共に、並行してスキャンを行うという形になる。
ここで、BgWorker側が nvcc でC/C++両方のホストコードを作ったり・・・などとしていると、下手すれば先にスキャンの方が終わってバックエンドを待たせてしまうというカッコ悪い事になるが、今回の NVRTC ライブラリを使うと、この辺の無用なオーバヘッドがない分、概ね1sec以下でビルドが終わってくれ、個人的には大変満足であった。
