スキャン速度10GB/sへの挑戦～その②～
一年半ほど前、次のようなエントリーを書いた。
かいつまんで言うと、多段JOINのように、実際に実行してみないと結果行数が明らかではなく、かつ、ステップ(k+1)の最適な問題サイズがステップkの結果に依存する場合、Kepler以降のGPUでサポートされた Dynamic Parallelism を使えば素直に実装できるという話である。
実際、この時期以降のGpuJoinロジックはDynamic Parallelismを用いて実装されていた。
だが、プロファイラ等を用いて詳しく調べてみると、どうやら、ある程度以上のパフォーマンスでクエリを処理している状況においては、このような設計はGPUが持つ本来の能力を引き出す上で必ずしも適切ではないという事が明らかになった。
例えば、以下のグラフは半年ほど前にStar Schema Benchmark(SSBM)を用いてクエリ処理のスループットを計測したものであるが、理論帯域2.2GB/sのSSD*1を2枚用いているにも関わらず、3.8GB/s辺りで処理性能が頭打ちになってしまっている。13種類のクエリを実行しているが、ひどいものだと1.5GB/s近辺にまでパフォーマンスが落ち込んでいる。
この時に使用していたGPUはTesla K80という少し古いもの*2なので、その辺は少し割り引いて考える必要がある。ただ、この程度の領域でサチっていては10GB/sのクエリ処理スループットなど"夢のまた夢"になってしまうので、数ヵ月前、この辺をプロファイラを用いて詳しく調べてみた。
一見すると、非常に細かな単位でびっしりとGPU Kernelが実行されており、GPU使用率が100%に張り付いているのも納得できるように思う。
Dynamic Parallelismを使ってGpuJoinを書き直したあたりもそういう認識で、あまり深掘りしてはいなかった。
しかし、実は妙なのである。
GPU時間を100%使用している事になっているgpupreagg_mainという関数は、基本的には、Dynamic Parallelismを用いて別のGPU kernelを起動するだけの"軽い"関数で、それぞれ4.7%、3.9%のGPU時間を使っているサブカーネルの方が実際のGROUP BY/集約演算の処理を行っている。
プロファイラで採取したタイムラインを拡大してみると、謎が解ける。
gpupreagg_mainがサブカーネルの実行完了を待っている間、この人がGPUで同時実行可能なGPU kernelをガメてしまい、その分、JOINやGROUP BYといったSQLワークロードを実際に処理するためのカーネルがまるで多重処理されていないのである。
なぜこれほどサブカーネルの同期処理に時間がかかってしまうのか。
Dynamic Parallelismを利用して起動したサブカーネルの完了をGPU内で待ち合わせる唯一の方法が、CUDAランタイム関数 cudaDeviceSynchronize() なのだが、この人は『Blocks until the device has completed all preceding requested tasks.』という動きをする。
つまり、無関係な並行タスクも含め、現在進行中の全てのタスクがいったん完了するのを待つ、という事になる。
ガツンと巨大なタスクをぶん回すHPC的なワークロードではあまり困る事もないだろうが、せいぜい数msで完了する細切れのGPU kernelを非同期かつ大量に発行するPG-Stromのワークロードでは、明らかに同期ポイントまでの遅延時間が長くなり、結果として、無駄にGPUを占有する時間も長くなってしまうように思える。
という事があり、ここ２ヵ月ほど大規模なPG-Stromのリファクタリングに取り組んでいた。
Dynamic Parallelismベースで実装した過去のGpuJoinやGpuPreAggは一旦お破算とし、GPUカーネル内で同期ポイントを持つ事なくJOINやGROUP BYのワークロードを実行するような実装へと作り替えた。
（先日のエントリ『GpuJoinの結果バッファ問題を考える。 - KaiGaiの俺メモ』などはその要素技術。）
で、再度改めてSSBMでベンチマークを行ってみた結果が以下の通り。じゃん。
システム構成は以下の通り。同じく理論帯域2.2GB/sのSSDを3枚搭載しているため、合計で6.6GB/sまでは読出しの帯域を持っている事になる。
- Server: Supermicro 5018GR-T
- CPU: Xeon E5-2650v4 (2.2GB, 12C) x1
- RAM: 128GB (16GB DDR4-2133, ECC; x8)
- GPU: NVIDIA Tesla P40 (3840C, 24GB) x1
- SSD: Intel SSD 750 400GB x3
- OS: CentOS 7.3 (kernel: 3.10.0-693.2.2)
- SW: CUDA 8.0 (driver: 384.66)
- DB: PostgreSQL 9.6.5, PG-Strom 2.0devel
それぞれ合計で3回ずつクエリを実行し、その中央値を結果として採用した。
スループットは以下のように導出している『(lineorderテーブルサイズ[MB])/(クエリ実行時間[sec])』
SSD-to-GPUダイレクトSQL実行を使った場合、物理限界6.6GB/sに対して、おおむね6.0GB/s前後のクエリ処理スループットを発揮できている。また、半年前の計測ではサチっていたQ4_1やQ4_2といったクエリでも他と変わらないパフォーマンスを出せている。
一方で、ノーマルのPostgreSQLの場合は2.0GB/s前後の処理スループットに落ち着いている。
今回は、手持ち機材の関係で物理限界 6.6GB/s までの計測となったが、このCPU自体は9.5GB/s～9.7GB/s近辺まではP2P DMAでデータ転送できることが分かっているので、新しいSSDを入手する事ができれば、処理スループット上限の更新に取り組んでみたい。
このベンチマーク実行中のGPU使用率は概ね35%～40%程度。まだ余裕はある。
