Dynamic ParallelismでGpuJoinはどう変わるか

NVIDIA GPUのKepler/Maxwell世代で*1対応となったDynamic Parallelismという機能がある。

GPUバイス上で動作するプログラム(Kernel関数と呼ぶ)を起動する際には、そのKernel関数を実行するために同時に何個のスレッドを起動するかを指定する。
例えば、10万レコード分のデータをGPU側に転送して各レコードに対して条件句を評価したいという場合なら、スレッド数=10万を指定すれば、GPU側では実装されている数百~数千コアをタイムスライスで切り替えつつ10万個のスレッドを処理してくれる。

従来、Kernel関数を起動する事をできるのはCPU側でCUDA API関数を用いた場合に限られていたのだが、前述のKepler/Maxwell世代のGPUを使えば、GPU上で実行中のKernel関数から別のKernel関数を起動する事ができるようになる。

f:id:kaigai:20160206172259p:plain

事前に問題サイズが明らかな場合にはそれほど嬉しい機能ではないのだが、問題は、Kernel関数を複数のステップに分けて実行しなければならないケースでかつ、問題サイズが前ステップの結果に依存する、すなわち、事前に問題サイズを正確に予測する事が難しいケースで非常に役に立つ。

典型例が、複数段から成るJOIN。

f:id:kaigai:20160206172312p:plain

上記の例は、table-0~table-3まで4個のテーブルをJOINする場合の処理を模式的に記したもの。
最初のステップでtable-0とtable-1を結合するケースでは、table-0とtable-1の行数は事前に明らかなので、適切なスレッド数を指定した上でKernel関数を起動してやればよい。
しかし、次のステップでは、table-0とtable-1をJOINした結果とtable-2のJOINを行うため、実際に入力される行数は実行されるまで分からない。最後のステップでも、table-0とtable-1とtable-2をJOINした結果とtable-3のJOINを行うため、これも実際の入力値となる行数は事前には分からない。
一般的に、不確定要素の上に不確定要素を積み重ねていくほど誤差は大きくなり、クエリ実行計画の最適化というのもこういった誤差との戦いである。閑話休題

これに対処するには2通りの方法が考えられる。

  1. 投機的実行:予め問題サイズを予測しておき、仮にダメなら一度失敗した時の情報を利用して再実行する。
  2. 同期実行:各ステップを完了する毎にCPU側でこれを捕捉し、前ステップの情報を利用して次ステップを起動する。

というものだった。どちらにも一長一短がある。
①投機実行の場合、当たれば非常に高速に動作するが、スレッド数・バッファサイズなどが推定値から大幅に外れた場合、再実行により悲劇的な実行性能となってしまう。一般的に、安全側に倒してスレッド数・バッファサイズを多目に取っておくと、今度はリソースを馬鹿食いする。例えば、前ステップで1000行しか生成されないにも関わらず、100万スレッドを要するという予測でKernel関数を実行した場合、99万9000スレッドは起動されたものの何もせずに終了する。このコストも馬鹿にはできない。
一方、②同期実行の場合は各ステップ毎にCPU側へ結果を返し、それに基づいて次ステップのKernel関数を起動するので、まぁ、リソース量の予測という点では一番良いのだが、肝心の性能がお話にならない。

Dynamic Parallelismが加わる事で、前ステップの結果を確認し、それに基づいて次ステップの実行に必要なスレッド数を計算*2した上で、次ステップのKernel関数を起動すればよい。

さて、ではDynamic Parallelismバンザイかというというと、ちょっと特殊な使い方をする必要がある。
なぜなら、CUDAコードの実行時コンパイルをサポートするNVRTC(NVIDIA Run-Time Compiler)ライブラリが次のように仰っておるので。

http://docs.nvidia.com/cuda/nvrtc/index.html#known-issues

6. Known Issues

The following CUDA C++ features are not yet implemented when compiling with NVRTC:

  • Dynamic parallelism (kernel launches from within device code).

が、普通にnvccで静的ビルドしたバイナリでDynamic Parallelが使えるのに、NVRTCでビルドしたバイナリで使えないのはおかしい・・・やり方があるハズ、と思って調べると、ちゃんとCUDA C Programming Guideに記述があった。

曰く、デバイス側のランタイム関数であるところの

extern "C" __device__ void *
cudaGetParameterBuffer(size_t alignment, size_t size);

で、Kernel関数起動用の引数バッファを獲得し、

extern "C" __device__ cudaError_t
cudaLaunchDevice(void *kernel_function,
                 void *parameterBuffer,
                 dim3 gridDimension,
                 dim3 blockDimension,
                 unsigned int sharedMemSize,
                 cudaStream_t stream);

を使って、別のKernel関数を起動する。

起動したKernel関数は非同期で実行される事になるが、同じくデバイス側のランタイム関数である

extern "C" __device__ cudaError_t
cudaDeviceSynchronize(void);

を使えば、非同期実行中のKernel関数を同期する事ができる。

これらの関数の実体は libcudadevrt.a ライブラリで提供されているので、従来、CUDAプログラムをNVRTCでコンパイル、生成されたPTX*3を cuModuleLoadData() 関数に与えてロードしていたパスが少し変わってくる。

f:id:kaigai:20160206172307p:plain

といっても、それほど大変な違いではなく、PTX形式のデータと外部のライブラリをリンクするために cuLinkXXXX() 関数を呼び出し、cubinというまた別の実行形式フォーマットを生成して、最終的に同じ cuModuleLoadData() 関数へロードする。

で、ここまで頑張った結果のパフォーマンスは以下のような結果となった。

f:id:kaigai:20160206172318p:plain

いつものように、JOIN+GROUP BYの性能をみるためだけのマイクロベンチマークを実行。

EXPLAIN ANALYZE
SELECT cat, avg(ax) FROM t0 NATURAL JOIN t1 [NATURAL JOIN ...] GROUP BY cat;

上記のクエリを、JOINするテーブルの数を増やしながら応答速度を計測してみた。
t0テーブルは1億件、その他のテーブルは10万件を含んでおり、ファクトテーブルとマスターテーブルの結合をイメージした処理。
比較対象は(青)PostgreSQL v9.5、(赤)PG-Strom Dynamic Parallelなし、(緑)PG-Strom Dynamic Parallelありで、縦軸は応答速度。つまり、短いほどベター。

その他の条件に関しては以下の通り。

グラフから見て分かるように、PG-Stromを使った場合、素のPostgreSQL v9.5と比べると、JOINするテーブルの数が増えたとしても、処理時間の増加は非常に緩やかな値となっている。
ただ、Dynamic Parallelismなしのケースでは、JOIN処理の段数が8を越えた辺りで問題サイズの推定に失敗し始め、リトライにより急速に処理時間が悪化するのが分かる。
一方、Dynamic Parallelismありのケースでは、少なくともこの程度の問題の複雑さではびくともしない事が読み取れる。

一番右側のt0...t8までを結合するケースでは、素のPostgreSQL v9.5 254.44secに対し、従来のDynamic ParallelなしのPG-Stromでは70.31sec、Dynamic ParallelありのPG-Stromでは18.81secで処理を完了している。

残念ながら良い事ばかりではない。結合段数が少ない場合、例えば2テーブルでのJOINなどは、Dynamic ParallelによりGPU側でのオーバーヘッドが増えているために、Dynamic Parallelなしの実装に比べ、多少、性能が劣化している。
また、実装をDynamic Parallelに切り替えるという事は、一部の古い Kepler 世代のGPUが非対応になる事を意味しており、現時点でインパクトが大きいのはAWSGPUインスタンス*4が使えなくなる事。ただ、これは時間が解決すると期待したい。

*1:正確に言うとComputing Capability3.5以降。つまりKepler世代でもTesla K20以上のハイエンドモデルとMaxwell世代の全モデルが対象

*2:ただし、流石にデバイスメモリの獲得は無理だ。バッファが足りなければ一部だけ実行し、CPU側へ残りの部分の再実行を促すという処理が必要になる

*3:NVIDIA GPU用のアセンブラ言語と思ってください

*4:GRID K520が載っており、CC3.0でDynamic Parallel非対応