GTC2017 - 個人的トピック

個人的トピックまとめ。

ポスター発表

今回、急遽渡米することを決めた理由がこのポスター発表。

セッションやトレーニングと並行して、GTCではGPUを用いた研究開発ネタのポスター展示が行われており、今年は約140件のポスター展示が採択されている。
このうち、プログラム委員の評価が高い20件が、S7480 - Fast Forward Poster Program for the Top 20 Posters のセッションで各々4minづつの発表を行う。さらに Top-20 の中からプログラム委員が事前に選定した5件が "Top-5 Poster Finalist" という扱いで、コンベンションセンターの正面、最も目立つ場所に掲載され、参加者がモバイルアプリ経由でどのポスターが気に入ったかを投票する。
最も票を集めたポスターには GTC2017 Poster Winner Award と副賞$5,000が贈られることになる。

今年は旅費の関係もあって元々渡米する予定はなかったので、セッション発表には申し込まず、ポスターを2枚投稿しただけだった。

このうち、SSD-to-GPUダイレクトSQL実行機能を紹介した P7130 - An intelligent storage for PostgreSQL database がTop-5に選定されたという連絡が来たのが4月29日(土)。Top-20の4分発表だけならちょっと考えたが、Poster Winner Awardの可能性があるならという事で急遽渡米準備。フライトとホテルを抑えてサンノゼに向かったワケである。

ただ、残念ながら結果は及ばず。
NVIDIA-JPの皆様からも他のゲストにご紹介頂いたり、偶然その場にいたNECのメンバに応援してもらったものの、7000人のGTC参加者の投票の結果*1、台湾の学生さんのMACHINE LEARNING領域の研究がPoster Winner Awardに選ばれた。おめでとうございます。
最終選考は一般からの投票だった訳だが、やはり応援してもらったのに結果につながらなかったというのは申し訳ない。次に機会があれば、今度はきちんとTop-1を取れるようにしたいものである。

ただ、今回ポスター発表という事で気付いた事も。GTC16, GTC15と過去2回の参加はいずれもセッション発表のみの参加で、これはある種一発勝負の要素がある。つまり、同じ時間帯に注目度の高いセッションが入っていたり、講演タイトルの"引き"がいま一つだと、わざわざ米国まで出かけて聴衆が10人くらいだとか、そういう悲しい事になる可能性も多いにある。

一方で、空き時間に人がぶらぶら眺めに来るポスターであれば、ちょっと熱心に読んでいる人がいたら『Hello, I'm author of the poster. Let me introduce our research briefly.』なんて声をかければ*2何回でもディスカッションのチャンスはあるし、場合によっては、セッションにただ参加して聞いているだけの人よりも高いレベルの認知/関心を得られるかもしれない。

来年以降、研究開発の発表を行う時には、単にセッションにプロポーザルを出すだけじゃなく、これまでリーチできなかった人の認知/関心を得るためにどういった発表の仕方がベストかもう少し考えてみた方がよさげ。例えば、ピュアな技術の部分はポスターとして、製品や事例に関してはセッションでというやり方があり得るかもしれない。

ちなみに、ポスターに申し込む時には一つ注意があり、プロポーザルを出す段階で既にポスターを作成していなければならないという事。
確か１月下旬が〆切だったが、担当のKaraさんから『ポスターのPDFが添付されてないですよ』と突っ込みがあり、ポスター2枚を大慌てで仕上げる羽目になった*3。ただ、これもちょっと考え物で、１月以降の進捗が発表に反映されないというデメリットはあるとは思う。

[idea] Visual Profiler

その他、参加したセッションの中でPG-Stromに応用が利きそうなものが二つ。
一つは Visual Profiler で、不勉強にして今まで知らなかったのだが、あるマシンでCUDAプログラムを実行し、そこでプロファイラが生成したログファイルを、別のマシンにインポートして可視化するという機能があるらしい。

確かに、CUDA Driver APIを見てみると、以下のようにプロファイラを初期化する関数の中に、出力ファイルを指定するオプションがある。

CUresult
cuProfilerInitialize(const char* configFile,
                     const char* outputFile,
                     CUoutput_mode outputMode);

今まで、PostgreSQL v9.5ベースの実装（つまりCPUパラレルなし）では、例えばデータ転送時間にしてもDMAの開始時刻/終了時刻を記録して、総転送データサイズを時間で割ってやれば、プロファイリングとして比較的直観に一致する結果が得られていた。
しかし、CPU+GPUハイブリッド並列に対応するPostgreSQL v9.6以降では、他のバックエンドとDMA転送がかち合った結果、PCIeバス使用率は高い（= 実際には10GB/sとか出る）一方で、処理時間が間延びした分転送レートが非常に低く見えてしまうという問題があった。
EXPLAIN ANALYZEで手軽に性能情報を確認できるというメリットがある一方で、CPU並列と相性が悪いというのは未解決問題であった。
であれば、細かい事は全部CUDAの実行系に任せて、PG-Stromの側でプロファイリングを頑張らないのも一つの選択肢ではないのかとも思えてきた。

実際、GPU kernel実行時間やDMA Send/Recvといった定量指標だけでなく、GPUタスクがうまい具合にキューに突っ込まれてSSDやGPUを遊ばせていないかというのも非常に重要なチューニングポイントではあるが、現状のPG-Stromのプロファイラはこれをきちんと取る事ができない。

[idea] GDF(GPU Data Flame) プロジェクト

MAP-DやH2Oといった、データ解析分野でGPUを使っている人たちが始めたプロジェクトで、要は、GPU上のグローバルメモリに配置したデータをGDFに対応したアプリケーション間で交換するという規格らしい。
確かに、cuMemAlloc()で獲得したメモリ領域は、cuIpcGetMemHandle()を使って識別キーをエクスポートする事が可能で、識別キーを受け取った他のプロセスはcuIpcOpenMemHandle()を呼び出す事でその領域を自分のプロセスでも利用する事ができるようになる。
彼らのアイデアとしては、GPU上の統計解析・機械学習アプリケーションを連携させるために、毎度毎度GPU RAM⇔CPU RAM間のデータ交換は避けたいというもの。GDFの仕様に従ってデータ領域をエクスポートすれば、連携対象のアプリケーションも同じようにその領域を参照できるはずだというもの。

海外としては二つ狙いがある。
一つは、やはり開発リソース・得意分野の関係で自分たちが統計解析・機械学習のエンジンを全て自作するか？？というもので、サードパーティと連携できる芽があり、さらにNVIDIAもこれを推しているなら悪くはないだろうという判断。
もう一つは、PL/CUDA関数で1GBを越えるデータサイズを扱う時に問題となっている、PostgreSQL可変長データサイズの問題を実質的にクリアできる可能性があること。

例えば、10GB程度の行列をGPUにロードすることを考える。Tesla P40であれば24GBのデバイスメモリを搭載しており、メモリの割り当て自体はそれほど大変な話ではない。しかし、PostgreSQL可変長データのサイズ制限により、これをPL/CUDAの引数として渡そうとすると、実質的には1個の行列であるにも関わらず複数の2次元配列に分割しなければならないなどの制約が生じる。
では考え方を変えて、GPUデバイスメモリ上に確保したメモリ領域を、PostgreSQLからForeign Tableとして読み書きできるようにしてはどうか？
この方法であれば、一度に受け渡すデータサイズが1GBを越える必然性はないし、FDWドライバが内部的に管理するデータ構造であれば(varlena型ではないので)1GBの壁は生じない。

GPU RAM上にデータを保持するForeign Tableを使用する場合、おそらく、PL/CUDA関数への引数はregclass型を用いてForeign TableのOIDが渡るようにすべきだろう。元々データはGPU RAM上に存在していることが前提なので、これらのデータは毎回RAM⇔GPU間を転送する必要がなくなる。
残念ながらTop-20に残らなかった方のポスター（P7129）では、PL/CUDAを用いた類似化合物検索を行っているが、1GB制約に合わせるため化合物データ1000万件(1.5GB)を4分割してPL/CUDA関数に与えている。しかし、予め化合物データがGPU側にロードされていれば、呼び出しのたびに2次元配列をセットアップする必要はなくなり、大幅に統計解析処理の効率が上がるだろう。

ここ最近の一連の開発で学んだのは、統計解析・機械学習のワークロードにおいては何がしかの前処理は必ず必要になるため、GPU RAMと関連付けられたForeign Tableに対して SELECT ... INTO fdw_table FROM transaction_table するというのは、どうやらユーザの負担を大きく増やすという性質のものではなさそうだという点。
生データからサマリを作ったり、前処理を行うというのはPG-Stromの得意とするところなので、これをPL/CUDAと組み合わせるための良い閃きが得られたと思っている。

[idea] 機械学習でのSSD-to-GPUダイレクトの利用

ポスターの発表に関連して、『これは機械学習のデータロードに使えませんかね？』と、以前にSSD-to-GPUダイレクトをブログで書いた時に頂いたのと全く同じ反応を、しかも複数の方から頂いた。
時間が取れないこともありペンディングになっているが、それほど難しい話ではないので、もう一度真面目に検討してみないと…。

[idea] CUDA MPSの利用

Tesla V100とCUDA9.0の発表に関連して、MPSの機能拡張によって、SMに空きがあるときはどんどんGPU kernelを放り込んでいけるようになるという説明があった。
現状、PostgreSQL v9.6に対応したバージョンはMPSと同じような機能を実現するバックグラウンドワーカを使っているが、こういった機能強化を取り込むためにも、できるだけ標準的なソフトウェアスタックの上に構築したいという思いはある。
現状、MPSだとストリームにコールバックを突っ込む事ができない。ここを何とかカバーする方法があれば、、、、という事で要検討。

*1:投票率は知らないがｗ

*2:なお、話しかけた相手が日本人だと若干気まずい。

*3:自業自得である

GTC2017 - Tesla V100 と CUDA9.0

今年のGTCはジェンスン・ファンCEOの基調講演が3日目に設定されており、そこでVolta世代の新製品 Tesla V100 と、それを搭載するDGX-1などサーバ製品とGPUクラウドが発表された。
4日間の日程のうち、情報の解禁が3日目の正午なので、その後は慌ただしくVoltaのアーキテクチャや、関連するCUDA9新機能のセッションに参加。
聞くべきセッションが被っちゃったりもしたので、できれば基調講演は初日の方が日程に余裕はあったかなぁ…。

Tesla V100について

5120CUDAコア、16GB(HBM2; 900GB/s)のメモリを搭載するまさに化け物プロセッサ。

デバイス全体で80個のSMを搭載する。SMあたりの構成は64CUDA CoreなのでPascalと同等だが、目新しい基軸としてTensor Coreという新しい計算ユニットが追加されている（SMあたり8個 = 全体で640個）。また、L1 Cacheが128KB x 80 = 10MB、L2 Cacheが6MB搭載されている。

これは4x4行列同士のFMAD演算(←と言っていいの？)を実行する専用のH/Wユニットで、従来であれば、16個のCUDAコアが協調して掛け算/足し算を行うべき箇所を命令一発で実行するというもの。

深層学習のアルゴリズムに関して詳しくないんだけども、ハマる状況であれば綺麗に高速化できるんであろうし、実際、TensorCore機能を使った機械学習ワークロードの実行速度が、P100比で9.3倍に向上したという事はアナウンスされている。

また、6本のリンクを全てフル帯域で使った場合の理論帯域だけども、NVLinkで300GB/sのスループットがあり、NVLink接続のリモートGPUに対するcache coherencyやatomic operationが可能だとのこと。

この辺の発表を聞いていての印象は、15年-20年ほど前からx86 CPUが辿った道のりに似てるなぁ～と。

Intel CPUの場合、マルチメディア処理命令としてMMX/AVXという拡張命令が加わったが、それを自分のプログラムから直接使うという人はあまり多くなく、実態としてはglibcなどのライブラリが内部的にこれらの機能を使うようになったために、それを通じて、ハマるケース限定でハードウェア専用命令を使うようになってきている。
TensorCoreの場合も、cuDNNやcuBLASなどのライブラリ群が内部的にこの機能を使うとのことなので、ユーザプログラムが行列計算や深層学習にこれらのライブラリ群を使っていれば、知らず知らずのうちに恩恵を受けることになる。というか、ラップしてくれないとちょっと使える状況が特殊すぎて辛い。

もう一つ。NVLinkの帯域が向上し、しかもUnified Memoryでアドレス空間の問題が解決し、cache coherencyやatomic operationに対応するとなれば、他のGPUデバイス上のメモリをローカルメモリに近い感覚で使うことができるようになるという事だろう。
そうすると、かつてLinuxがNUMA環境への対応を行ったように、メモリの物理的な距離を意識しなければならない状況が生まれつつあるように思える。例えば Tesla V100 は 16GB(HBM2) のメモリを積んでいるが、アプリケーションが20GBのデータをロードして2台のGPUでぶん回したい場合、データ量は明らかに1台のGPUに収まらないので2台に分散ロードする必要がある。
この時、物理GPUページの割り当てが『前半10GBはGPU0、後半10GBはGPU1』みたいなざっくりで上手く動作するのだろうか？もしかしたら、メモリアクセスの局所性が期待できるタイプのワークロードなのか？　それとも、全体を舐め回すのでインタリーブ型の物理データ配置を行うべきなのか、それほど簡単な問いではないように思える。


SMの構成は図の通りであるが、L1 Cacheの扱いが少し変わっている。
PascalではL1 Cacheが24KBだったのが、128KBあるL1/Shared Memory共通の領域のうち、共有メモリが使用していない領域をL1 Cacheとして利用できるわけである。
Maxwell以前にL1/Shared Memoryの大きさを16KB/48KB、32KB/32KB、48KB/16KBの中から選んでいたのに通ずるが、もう少しフレキシブルに"早いメモリ"を使ってくれているようである。
ただ、GPUのL1キャッシュの場合(少なくともPascalまでは)、キャッシュの下に転がっているGlobal Memoryの内容がGPU kernel実行中に少なくとも変更されない/Read-Onlyであるという制約があったが、これはおそらくVoltaでも変わらないものと思われる（何も言及無かったし）

また、Pascal以前のGPUと大きく変わってくる点として、プログラムカウンタ(PC)がwarp単位ではなくthread単位になるという違いがある。
従来は、if (...) { ... } else { ... } ブロックを実行するときに、warp中の一部スレッドだけが条件が真、他のスレッドが偽である場合には両方のブロックを逐次的に実行することになるため、他方のスレッド群が遊んでしまい、それ故にGPUは条件分岐に弱いと言われてきたが、各スレッドがPCを持つという事はこうした弱点に手を打ってきたという事でもある。
（とは言え、TrueブロックとFalseブロックの分量があまりにアンバランスだとアカンでしょうが）

5月13日追記：
FBでNVIDIA-JPの成瀬さんから教えていただきました。
『プログラムカウンタ(PC)がwarp単位ではなくthread単位になる』というのは必ずしも間違ってはいないのですが、CPUスレッドのように各CUDAコアが全く独立な命令を実行できるという事ではないようです。

より正確に述べると、PC(ここでは次に実行する命令のアドレスとしましょう)は、Voltaではスレッド毎に持ちます。ただし、同じwarp内のスレッドは、PCが同じであるスレッドだけが、同じクロックで実行されます。warp内で、別PCの命令が実行されることは、これまで通り、ありません。

つまり、if (...) {} else {} ブロックで分岐が発生する場合、従来、trueブロックを実行中のCUDAコアがある時はfalseブロックに分岐したCUDAコアは寝ているしかなかった。Voltaにおいては、例えばDRAMアクセスなどでメモリストールが発生しておりtrueブロックを実行中のCUDAコアが寝ている状況では、その間を利用してfalseブロックに分岐したCUDAコアも先に実行を進めることができる、という事のようです。

他にニッチなところとしては、SMが80個にも増えてアプリケーションによっては全てのSMを使いきれないという事を想定しているのか、MPS(Multi Process Service)を使用した時に複数のアプリケーションからのGPU kernel起動要求を調停して、GPU資源を切り分けるという使い方ができますよーとの事。ただ、MPSはcuStreamAddCallback()のようにコールバック関数を呼び出す事ができなかったハズなので、非同期処理が中心のPG-Stromとは今一つ相性がなぁ…。

CUDA 9について

Voltaアーキテクチャ向けのTensorCoreを使うための機能や、cuBLASなどライブラリのTensorCore対応が行われるらしい。

ただ、最も革命的と言ってよい変化は、今までスレッド間の同期は __syncthreads() を使ったSM内の同期のみであったのが、SMを越えて複数のグリッド間で／複数のGPUデバイスで並行動作するカーネル間で同期を取ることが可能になったというもの。

GPU kernelを起動するときは、必要なスレッド数がN個としてN=(ブロックサイズ)×(グリッドサイズ)
という関係になるようスレッド数を与えて起動する。ブロックサイズの上限は 1024 であり、GPU kernelの実行時にはブロックサイズに依らず、特定のブロックは特定のSMで実行される。
（グリッドサイズ > デバイスのSM数 な場合は、ブロックを順番に実行する）

で、大雑把に言えば、従来はH/Wに由来するブロックを単位とする同期の他には、GPU kernel全体の実行完了を待つという同期方法しか取れなかったが、CUDA 9で導入される co-operative group という概念を使えば、H/W的な構造と論理的な処理の同期の単位を切り分けることができる。

これがなぜ嬉しいかと言うと、例えばBitonic-Sortingを実行するときに、要素数が1個のSMで処理するには辛い数（1万個とか）になってくると、ステップ事に一度GPU kernelを完了し、そこで同期した上で改めて次のステップを実行するという処理が必要になる。
これは結構しんどい処理で、PG-StromでGPUでのORDER BY処理の実装を断念した要因の一つでもある。
が、SMを越えた単位で処理を同期できるとなれば、ステップ事に同期ポイントを設定してやればよいだけである。

例えば、以下のような感じで処理を記述する。

__global__ void
particleSim(Particle *p, int N)
{
  grid_group g = this_grid();

  // 1. Jobs prior to sync
  for (i = g.thread_rank(); i < N; i += g.size())
    integrate(p[i]);

  g.sync() // 2. Sync whole grid

  // 3. Jobs post sync
  for (i = g.thread_rank(); i < N; i += g.size())
    collide(p[i], p, N);
}

grid_groupというのはco-operative groupを表現するためのクラスで、this_grid()はこのGPU kernelを起動した時の全てのスレッドを含むオブジェクト。
つまり、g.sync()は複数のSMに跨った全てのスレッドで同期を取ることを意味している。

他にも、従来同様に現在のブロックに含まれるスレッドを表現する thread_group や、逆に複数のデバイスで同時に実行中の各GPU kernelに含まれるスレッドを表現する multi_grid_group が定義されている。
なお、SM間/デバイス間での同期機構を使う場合には、GPU kernelを起動する時に用いるAPIが従来のものとは少し異なっており、Runtime APIでは以下のAPIを用いてGPU kernelをキックしなければならない。

  • grid_groupを使う場合
    • cudaLaunchCooperativeKernel()
  • multi_grid_groupを使う場合
    • cudaLaunchCooperativeKernelMultiDevice()

で、更に面白いことに、このco-operative groupというのは更に細かい単位にパーティション化する事ができる*1。

こんな調子で、thread_group gに含まれるスレッドを tiled_partition() を介して4つのパーティションに分割する事ができるらしい。
なので、真ん中の tile.sync() ではパーティション毎の同期という事になる。

  thread_group g = this_thread_label()
  thread_gruup tile = tiled_partition(g, 4);
        :
  tile.sync();
        :

現状、CUDA 9では等分割のパーティショニングのみだが、将来的には labeled partition という形にコンセプトが拡張されるとの事。
"labeled" というのは、以下のように partition() 関数に与えた整数値によってそのスレッドが属するパーティションを決めるというもの。

  int label = foo() % 4;
  thread_group g = partition(this_thread_block(), label);

何が嬉しいかというと、GROUP BY/集約演算をGPUで実装するときに、GROUP BYで指定したカラムの値が同一のスレッド間でのみ同期を取り、縮約演算を行うというロジックが可能になること。これは、現状atomic演算を用いて実装せざるを得ないことを考えると非常に大きなアドバンテージになる。

CUDA 9.0の登場が待ち遠しいところである。

余談

SM間で同期を取るとき、GPU kernelのブロック数がデバイスのSM数よりも多い場合、ハードウェア的にはSMを明け渡さないといけないはずなので、そのタイミングで共有メモリやレジスタをグローバルメモリに退避する（= そんなに軽い処理じゃないのでは？）と思い、セッション後に発表者を捕まえて質問してみた。
曰く、Pascal世代で(H/W的には)対応したプリエンプションの仕組みを使って実装されており、レジスタや共有メモリの使用量がH/Wリソースを超える場合には確かにそうなる。なので、GPU kernelを起動するときのパラメータチューニングが必要で、SMあたり {（ブロック数）/（デバイスのSM数）} のブロックを保持できる程度のレジスタ/共有メモリの消費に抑えておく必要がある。
一応、その辺は Occupancy を計算するAPIを使って効率的なブロックサイズ/グリッドサイズを求めてほしいとの事。

お願い

『S7285 - Unified Memory on the Volta architecture』のセッションが、CUDA9のセッションと被ってて聞けなかった！誰か内容教えて！！

NVIDIA TechBlogにも既に発表内容を踏まえた記事が掲載されています。

(*) 本文中の画像は全て上記ブログからの引用です。

*1:が、セッションの中ではthread_groupの例しか示されてなかったので、grid_groupやmulti_grid_groupでもできるかは要確認

Posters flying on GTC2017

I'm on GTC(GPU Technology Conference) this year again.

I didn't plan to trip because of travel budget, so only posters are submitted and approved.
(Posters are shown regardless of the attendance of the presenter.)

But I got an e-mail below...

Hi Kohei,

After extensive review from multiple committees and nearly 140 posters, your poster has been chosen as a Top 5 Finalist. Congratulations!!

Eventually, I arranged a trip in hurry, then now in San Jose.

The top-5 posters are shown at most conspicuous place of the convention center.

Participants can vote one of their favor poster, then the best poster this year is chosen through the democratic process. :-)

I and other top-20 presenters had a short talk on evening of the 1st day of the conference.
4-minutes are given to each presenter, and we introduce their research in this short time-slot.
(It is shorter than usual lightning talks!)

Our challenge tries to de-define the role of GPU. People usually considers GPU is a powerful tool to tackle computing intensive workloads. It is right, and widespread of workloads utilizes GPUs well like deep learning, simulations and so on.
On the other hands, I/O is one major workloads for database; including data analytics which try to handle massive amount of data.


We designed and implemented a Linux kernel module that intermediates peer-to-peer data transfer between NVMe-SSD and GPU, on top of GPUDirect RDMA feature. It allows to load PostgreSQL's data blocks on NVMe-SSDs to GPU's RAM directly.
Once data blocks get loaded to GPU's RAM, we already have SQL execution engine on GPU devices.
It means we can process earlier stage of SQL execution prior to data loading onto the system main memory. The earlier stage includes WHERE-clause, JOIN and GROUP BY. They likely work to reduce amount of data.
For example, when an WHERE-clause filters out 90% of rows (it is not a strange assumption in reporting queries) on scan of 100GB tables, we can say 90GB of I/O bandwidth and dist cache to be wipe out are consumed by junk data.
Once data blocks are loaded to GPU, it can filter out unnecessary rows (and processes JOIN and GROUP BY), then write back the results of this pre-process to host memory. The data blocks written back contains all valid data because unnecessary data are already dropped off.


In the results, it performs nearly wired speed of NVMe-SSD.
In case of single SSD benchmark, query execution throughput by vanilla PostgreSQL is around 700MB/s. On the other hands, SSD-to-GPU Direct SQL-execution runs same query with 2.2GB/s; which is almost identical with catalog spec of Intel SSD 750 (SeqRead: 2200MB/s).
In case of multi SSDs cases, it also increases the performance, however, deeper investigation *1 and improvement for md0-raid support are needed.


We intend to position this technology to make summary and pre-process of raw data/logs generated by sensor devices, mobile devices and so on. Once logs get summarized, we can apply statistical-analytics / machine-learning algorithms.
This type of workloads are advantaged with PL/CUDA which is one other our feature to utilize GPUs in-database analytics.
Eventually, GPU-accelerated PostgreSQL will be capable to handle entire life-cycle of data utilization (generation - gathering - summarizing/pre-process - analytics/machine-learning - visualization).
It makes system configuration much simple, reduce cost for production and operation, and deliver the power of analytics to wider and wider people.

My poster is below (enlarge by click):

If you are favored with our development, don't forget to vote.

Enjoy your GTC2017!

(Comment to the above photo by my wife: it looks like a zombie.)

*1:At that time, we measured with Teska K80; that is a bit old.

進捗）SSD-to-GPU ダイレクトSQL実行機能

ここ暫くブログでまとめていなかった、SSD-to-GPUダイレクトSQL実行機能の進捗について。

この機能をかいつまんで言うと、NVMe-SSDに格納されているPostgreSQLのデータブロックをGPU RAMに直接転送し、そこでSQLのWHERE句/JOIN/GROUP BYを実行することで見かけ上のI/O量を削減するという代物である。
NVIDIAのTesla/Quadro GPUが対応するGPUDirect RDMA機能を使い、SSD<=>GPU間のデータ転送を仲介するLinux kernel moduleを使えば、CPU/RAMにデータをロードする前にGPU上での処理を行うことができる。

しばらく前からScan系の処理には対応していたが、JOIN/GROUP BYへの対応を加え、さらにPostgreSQL v9.6のCPU並列にも追従したということで、簡単なベンチマークなら取れるくらいまで開発は進んできている。

という事で、現時点での実力がどの程度なのか、手元の環境を使って測定してみることにした。
使用したのはDELL R730にTesla K80*1とIntel SSD 750(400GB)を2枚搭載したモデル。
GPUDirect RDMAの制約から、SSDとGPUが同一のCPU又はPCIeスイッチに接続されている必要があるため、共にCPU2配下のPCIeスロットに接続している。

ベンチマークに使用したのはTPC-Hを簡略化したStar Schema Benchmark(SSBM)と呼ばれるテストで、中核となるlineorderテーブルのサイズが明らかにRAMに収まり切らないようscaling factorを調整している。

SSBMでは何種類かクエリが定義されているが、基本的にはWHERE句/JOIN/GROUP BYという、集計クエリの典型的なものである。例えば以下のようなものである。

SELECT sum(lo_revenue), d_year, p_brand1
  FROM lineorder, date1, part, supplier
 WHERE lo_orderdate = d_datekey
   AND lo_partkey = p_partkey
   AND lo_suppkey = s_suppkey
   AND p_category = 'MFGR#12‘
   AND s_region = 'AMERICA‘
 GROUP BY d_year, p_brand1
 ORDER BY d_year, p_brand1;

それを実行してみた結果が以下の通り。じゃん！

青色がPostgreSQL(ファイルシステム経由I/O)での実行結果、橙色がPG-Stromでの実行結果で、それぞれ色の濃い方がSSD x1枚での実行結果、色の薄い方がSSD x2枚(md-raid0)での実行結果。

これを見る限り、SSD x1枚の場合、PG-Stromは概ね理論限界*2に達しているが、SSD x2枚のスループットである4.4GB/sにはまだ届いていない。ただ、これはGPUがTesla K80というやや古いモデルであった事や、GROUP BYでの集約演算がNumeric型の総和や平均など、Kepler世代のGPUには比較的辛い処理だという事は考慮しなければいけないだろう。（例えばQ4-1やQ4-2ではI/Oよりも集計処理にボトルネックがあるように見える）
一方PostgreSQLの場合はSSD x1枚で600-750MB/s程度、SSD x2枚で1.6GB/s程度のスループットなので、PG-Stromの場合はI/Oネックな集計処理において2.3～3.0倍程度の優位性があるという事になるだろう。

この辺、SSDとGPU、あるいはCPUのどこに律速要因があるのか追及するのはなかなか骨の折れる作業ではあるが、近々、最新モデルであるTesla P40で実行できるようになるので、最新世代のGPUではどの程度のスループットまで耐えられるのか、試してみたいところである。

本日の記事のより詳しい内容は、先日、BigData基盤研究会#7の方で喋らせて頂いた資料を公開しているので、そちらも併せて参照していただければ。

www.slideshare.net

*1:ちょっと古い

*2:Intel SSD 750(400GB)のSeqReadスペックは2200MB/s

PCIeスロット接続型NVMe-SSDまとめ（2017年4月時点）

SSD

PCIeスロット接続型のNVMe-SSDのスペック等、各社どうだったけな～と探すものの、案外まとまった情報がないので自分でまとめてみた。ブックマーク代わりです。
基本的に各社のWebに掲載されているカタログスペックを転記。手作業なので内容の正確さに関しては保証しかねます。（指摘頂ければ直します）

Intel社

コンシューマ向け

モデル Intel SSD 750
400GB(MLC)		 Intel SSD 750
800GB(MLC)		 Intel SSD 750
1.2TB(MLC)
形状 HHHL HHHL HHHL
PCIe PCIe 3.0 x4 PCIe 3.0 x4 PCIe 3.0 x4
容量 400GB 800GB 1.2TB
SeqRead 2100MB/s 2100MB/s 2400MB/s
SeqWrite 800MB/s 800MB/s 1200MB/s
RandRead 420kIOPS 420kIOPS 440kIOPS
RandWrite 210kIOPS 210IOPS 290kIOPS
リリース Q3'15 Q3'15 Q2'15

エンタープライズ向け

モデルl DC P3608
1.6TB(MLC)		 DC P3608
3.2TB(MLC)		 DC P3608
4.0TB(MLC)		 DC P3700
400GB(MLC)		 DC P3700
800GB(MLC)		 DC P3700
1.6TB(MLC)		 DC P3700
2.0TB(MLC)
形状 HHHL HHHL HHHL HHHL HHHL HHHL HHHL
PCIe PCIe 3.0 x8 PCIe 3.0 x8 PCIe 3.0 x8 PCIe 3.0 x4 PCIe 3.0 x4 PCIe 3.0 x4 PCIe 3.0 x4
容量 1.6TB 3.2TB 4.0TB 400GB 800GB 1.6TB 2.0TB
SeqRead 5000MB/s 4500MB/s 5000MB.s 2700MB/s 2800MB/s 2800MB/s 2800MB/s
SeqWrite 2000MB/s 2600MB/s 3000MB/s 1080MB/s 1900MB/s 1900MB/s 1900MB/s
RandRead 850kIOPS 850kIOPS 850kIOPS 450kIOPS 460kIOPS 450kIOPS 450kIOPS
RandWrite 150kIOPS 80kIOPS 50kIOPS 75kIOPS 90kIOPS 150kIOPS 175kIOPS
リリース Q3'15 Q3'15 Q3'15 Q2'14 Q2'14 Q2'14 Q2'14

※ Optane DCモデルに関してはARKにデータが上がってから更新予定。

Samsung社

エンタープライズ向け

モデル PM1725a
1.6TB(V-NAND)		 PM1725a
3.2TB(V-NAND)		 PM1725a
6.4TB(V-NAND)		 PM1725
3.2TB(V-NAND)		 PM1725
6.4TB(V-NAND)
形状 HHHL HHHL HHHL HHHL HHHL
PCIe PCIe 3.0 x8 PCIe 3.0 x8 PCIe 3.0 x8 PCIe 3.0 x8 PCIe 3.0 x8
容量 1.6TB 3.2TB 6.4TB 3.2TB 6.4TB
SeqRead 5840MB/s 6200MB/s 6300MB/s 5600MB/s 5600MB/s
SeqWrite 2100MB/s 2600MB/s 2600MB/s 1800MB/s 1800MB/s
RandRead 1000kIOPS 1000kIOPS 1000kIOPS 1000kIOPS 1000lIOPS
RandWrite 140kIOPS 180kIOPS 180kIOPS 130kIOPS 130kIOPS
リリース Q3'16 Q3'16 Q3'16 Q3'15 Q3'15

HGST社

エンタープライズ向け

モデル Ultrastar
SN260(1.6TB)		 Ultrastar
SN260(3.2TB)		 Ultrastar
SN260(6.4TB)		 Ultrastar
SN150(1.6TB)		 Ultrastar
SN150(3.2TB)
形状 HHHL HHHL HHHL HHHL HHHL
PCIe PCIe 3.0 x8 PCIe 3.0 x8 PCIe 3.0 x8 PCIe 3.0 x4 PCIe 3.0 x4
容量 1.6TB 3.2TB 6.4TB 1.6TB 3.2TB
SeqRead 6100MB/s 6100MB/s 6100MB/s 3000MB/s 3000MB/s
SeqWrite 2200MB/s 2200MB/s 2200MB/s 1600MB/s 1600MB/s
RandRead 1200kIOPS 1200kIOPS 1200kIOPS 743kIPOS 743kIPOS
RandWrite 200kIOPS 200kIOPS 200kIOPS 140kIPS 140kIOPS
リリース Q4'16 Q4'16 Q4'16 Q2'15 Q2'15

Seagate社

エンタープライズ向け

モデル Nytro XP7200
7.7TB(MLC)		 Nytro XP6500
1.5TB		 Nytro XP6500
4.0TB		 Nytro XP7102
800GB		 Nytro XP7102
1.6TB
形状 FHHL FHHL FHHL HHHL HHHL
PCIe PCIe 3.0 x16 PCIe 3.0 x8 PCIe 3.0 x8 PCIe 3.0 x4 PCIe 3.0 x4
容量 7.7TB 1.3TB 3.4TB 800GB 1.6TB
SeqRead 10000MB/s 4.0GB/s 4.0GB/s 2500MB/s 2500MB/s
SeqWrite 2300MB/s 1.5GB/s 2.0GB/s 850MB/s 900MB/s
RandRead 950kIOPS 300kIPOS 275kIOPS 245kIOPS 245kIOPS
RandWrite 64kIPS 100kIOPS 75kIOPS 35kIOPS 40kIOPS
リリース Q3'16 ??? ??? Q3'16 Q3'16

ショックだ・・・Nytro XP7200の3.8TB版がサイトから消えている・・・。 orz

OCZ社

コンシューマ向け

モデル OCZ RD400a
128GB(MLC)		 OCZ RD400a
256GB(MLC)		 OCZ RD400a
512GB(MLC)		 OCZ RD400a
1.0TB(MLC)
形状 HHHL HHHL HHHL HHHL
PCIe PCIe 3.0 x4 PCIe 3.0 x4 PCIe 3.0 x4 PCIe 3.0 x4
容量 128GB 256GB 512GB 1.0TB
SeqRead 2200MB/s 2600MB/s 2600MB/s 2600MB/s
SeqWrite 620MB/s 1150MB/s 1600MB/s 1550MB/s
RandRead 170kIPS 210kIOPS 190kIOPS 210kIOPS
RandWrite 110kIOPS 140kIPS 120kIOPS 130kIOPS
リリース Q2'16 Q2'16 Q2'16 Q2'16

Plextor社

コンシューマ向け

モデル M8Pe(Y)
128GB		 M8Pe(Y)
256GB		 M8Pe(Y)
512GB		 M8Pe(Y)
1.0TB
形状 HHHL HHHL HHHL HHHL
PCIe PCIe 3.0 x4 PCIe 3.0 x4 PCIe 3.0 x4 PCIe 3.0 x4
容量 128GB 256GB 512GB 1.0TB
SeqRead 1600MB/s 2000MB/s 2300MB/s 2500MB/s
SeqWrite 500MB/s 900MB/s 1300MB/s 1400MB/s
RandRead 120kIPS 210kIOPS 260kIPS 280kIOPS
RandWrite 130kIOPS 230kIPS 250kIOPS 240kIPS
リリース Q4'16 Q4'16 Q4'16 Q4'16

Kingstone社

エンタープライズ向け

モデル DCP1000
800GB		 DCP1000
1.6TB		 DCP1000
3.2TB
形状 HHHL HHHL HHHL
PCIe PCIe 3.0 x8 PCIe 3.0 x8 PCIe 3.0 x8
容量 800GB 1.6TB 3.2TB
SeqRead 6800MB/s 6800MB/s 6800MB/s
SeqWrite 5000MB/s 6000MB/s 6000MB/s
RandRead 900kIOPS 1100kIPS 1000kIOPS
RandWrite 145kIOPS 200kIOPS 180kIPS
リリース Q1'17 Q1'17 Q1'17

※なぜかメーカーのサイトから消えている・・・。

AWSのP2.*インスタンスで PG-Strom を試す

GPU PostgreSQL

従前、AWSの提供するGPUインスタンス g2.* に搭載されているGPUはGRID K520というちょっと古いモデルで、PG-Stromは非対応だった。
理由は、一年ほど前にComputing Capability 3.5以降で対応のDynamic Parallelism機能を使うように全面的に作り直したからで、詳細は以下のエントリを参照。

kaigai.hatenablog.com

その後、昨年の10月にAWSは新世代*1のGPUインスタンスを新たにリリースした。

japan.zdnet.com

これでPG-Stromの動作要件を満たすようになった上に、特にメモリ搭載量で相応の強化が行われたため、例えばPGconf.ASIAで発表を行った創薬領域の類似度サーチのような、I/Oが支配的でないようなワークロードであれば相応の効果が見込める、ハズである。

発表から少し間が空いてしまったが、p2.*インスタンスで動作検証を行い、当該環境をAMIイメージとして公開してみた。

環境構築

基本的にはAMIを探してポチポチするだけである。


最初に、リージョンを p2.* インスタンス提供済みのリージョンに切り替える。
現状、東海岸の「バージニア北部」、西海岸の「オレゴン」、欧州の「アイルランド」での提供が開始されている模様。東京はまだである*2。
とりあえず、今回は「オレゴン」リージョンを使用。

また、AWSアカウントを作成直後は p2.* インスタンスの起動制限数が 0 に設定されているため、p2.xlargeインスタンスを起動できるよう、EC2ダッシュボードの『制限』から『制限緩和のリクエスト』を行う必要がある。豆知識。


次に、コミュニティAMIから『PGStrom』で検索すれば、夜なべして作ったAMIイメージがヒットする。
これは PG-Strom v1.0 + PostgreSQL v9.5.5 + CUDA 7.5環境を CentOS7(x86_64) 上に構築したもので、テスト用のデータも既にセットアップ済みのものである。


次に、起動すべきインスタンスを選択する。
p2.*インスタンスは以下の3種類提供されており、今回は『p2.xlarge』タイプを使用した。
なお、Amazon様も『高性能データベース』用途ときちんと書いていてくれておりますｗ

一応、PG-StromはマルチGPUにも対応しているが、CPU並列を使えないPostgreSQL v9.5系列ではGPUへのデータ供給の方が先にボトルネックになってしまうので、あまり意味はない。
この辺は、現在開発中の PostgreSQL v9.6 ベースの実装でうまくハンドリングできるようになる。

インスタンスの構成を確認して起動。ストレージはroot区画に24GBをアタッチするだけのシンプルなものがデフォルト構成。

しばらく待っているとインスタンスが立ち上がってくる。

『ec2-user』でログインする事ができる。

動作確認

早速動作確認を行う。ec2-userはlocalhost経由でPostgreSQLサーバに接続できるよう設定されている。

[ec2-user@ip-172-31-24-196 ~]$ psql postgres
psql (9.5.5)
Type "help" for help.

postgres=#

Tesla K80 GPUが搭載されている事が分かる。
物理的には一枚のカード上にGPU 2個を搭載するデバイスなので、一枚のカードを他の誰かとシェアしている事になる。

postgres=# SELECT * FROM pgstrom_device_info();
 id |                       property                        |    value
----+-------------------------------------------------------+--------------
  0 | Device name                                           | Tesla K80
  0 | Total global memory size                              | 11519 MBytes
  0 | max threads per block                                 | 1024
  0 | Maximum block dimension X                             | 1024
  0 | Maximum block dimension Y                             | 1024
  0 | Maximum block dimension Z                             | 64
  0 | Maximum grid dimension X                              | 2147483647
  0 | Maximum grid dimension Y                              | 65535
  0 | Maximum grid dimension Z                              | 65535
  0 | Maximum shared memory available per block             | 48 KBytes
  0 | Memory available on device for __constant__           | 64 KBytes
  0 | Warp size in threads                                  | 32
  0 | Maximum pitch in bytes allowed by memory copies       | 2147483647
  0 | Maximum number of 32bit registers available per block | 65536
  0 | Typical clock frequency in kilohertz                  | 823 MHz
  0 | Alignment requirement for textures                    | 512
  0 | Number of multiprocessors on device                   | 13
  0 | Has kernel execution timeout                          | False
  0 | Integrated with host memory                           | False
  0 | Host memory can be mapped to CUDA address space       | True
  0 | Compute mode                                          | Default
  0 | Alignment requirement for surfaces                    | 512
  0 | Multiple concurrent kernel support                    | True
  0 | Device has ECC support enabled                        | True
  0 | PCI bus ID of the device                              | 0
  0 | PCI device ID of the device                           | 30
  0 | Device is using TCC driver model                      | False
  0 | Peak memory clock frequency                           | 2505 MHz
  0 | Global memory bus width                               | 384 bits
  0 | Size of L2 cache in bytes                             | 1536 KBytes
  0 | Maximum threads per multiprocessor                    | 2048
  0 | Number of asynchronous engines                        | 2
  0 | Device shares unified address space                   | True
  0 | PCI domain ID of the device                           | 0
  0 | Major compute capability version number               | 3
  0 | Minor compute capability version number               | 7
  0 | Device supports stream priorities                     | True
  0 | Device supports caching globals in L1                 | True
  0 | Device supports caching locals in L1                  | True
  0 | Maximum shared memory per multiprocessor              | 112 KBytes
  0 | Maximum number of 32bit registers per multiprocessor  | 131072
  0 | Device can allocate managed memory on this system     | True
  0 | Device is on a multi-GPU board                        | False
  0 | Unique id if device is on a multi-GPU board           | 0
(44 rows)

早速、シンプルな集計のクエリを叩いてみるが、その前にメインのテーブルに対してpg_prewarmを実行しておくが吉。
gp2ストレージタイプの場合、バッファに載っていないデータをロードするのに割と時間がかかるようなので、I/Oネックになってしまうと、CPUとかGPUとかいう次元の遥か手前で引っかかってしまう。

postgres=# SELECT pg_prewarm('t0'::regclass);
 pg_prewarm
------------
     833334
(1 row)

実行計画を確認する。JOIN+GROUP BYをGPUで実行するようにプランが選択された事がわかる。

postgres=# EXPLAIN SELECT cat,count(*),avg(ax) FROM t0 NATURAL JOIN t1 GROUP BY cat;
                                         QUERY PLAN
---------------------------------------------------------------------------------------------
 HashAggregate  (cost=3578311.16..3578311.48 rows=26 width=12)
   Group Key: t0.cat
   ->  Custom Scan (GpuPreAgg)  (cost=14139.24..2873626.84 rows=234 width=48)
         Reduction: Local + Global
         GPU Projection: cat, ax
         ->  Custom Scan (GpuJoin) on t0  (cost=10139.24..2838812.04 rows=98599882 width=12)
               GPU Projection: t0.cat, t1.ax
               Depth 1: GpuHashJoin, HashKeys: (t0.aid)
                        JoinQuals: (t0.aid = t1.aid)
                        Nrows (in/out: 98.60%), KDS-Hash (size: 13.47MB, nbatches: 1)
               ->  Seq Scan on t1  (cost=0.00..1935.00 rows=100000 width=12)
(11 rows)

1億行のテーブルのJOINとGROUP BYで11.3秒要している。
セッションに接続後の初回実行だと、これに加えて、GPUコンテキストの初期化やGPUコードのJITコンパイルの時間も余分にかかるが、オンプレに比べるとAWSの環境ではこの辺が多少もっさりする印象がある。

postgres=# SELECT cat,count(*),avg(ax) FROM t0 NATURAL JOIN t1 GROUP BY cat;
 cat |  count  |       avg
-----+---------+------------------
 nnn | 3845736 | 49.9122204644341
 ccc | 3842975 | 49.9253161146813
 ddd | 3841209 | 49.9346989307005
 aaa | 3848221 | 49.9265219996308
 kkk | 3843481 | 49.9232348058601
 fff | 3845484 | 49.9434949581969
 iii | 3846743 | 49.9227199719054
 jjj | 3846076 | 49.9292863471083
 qqq | 3845646 |  49.945213365697
 hhh | 3842519 | 49.9266693322143
 ttt | 3846725 | 49.9232478784252
 ooo | 3847927 | 49.9346102129999
 zzz | 3847116 | 49.9320751724648
 lll | 3848447 |  49.927865906661
 www | 3846691 | 49.9537192102014
 bbb | 3845315 | 49.9281298849625
 ppp | 3850842 | 49.9149069792416
 eee | 3846285 |  49.931458202446
 xxx | 3845570 | 49.9577920754281
 ggg | 3845044 | 49.9409383291351
 rrr | 3847816 | 49.9341189910578
 uuu | 3845813 | 49.9295202543591
 vvv | 3849157 | 49.9253944163053
 yyy | 3843414 | 49.9364087656463
 mmm | 3848758 | 49.9033622681507
 sss | 3846990 | 49.9213589517191
(26 rows)

Time: 11326.834 ms

一方、PG-Stromを無効にした場合は以下のように、同じクエリの実行に51.0秒を要している。ﾄｩｯﾄｩﾙｰ。

postgres=# EXPLAIN SELECT cat,count(*),avg(ax) FROM t0 NATURAL JOIN t1 GROUP BY cat;
                                 QUERY PLAN
-----------------------------------------------------------------------------
 HashAggregate  (cost=3937016.94..3937017.26 rows=26 width=12)
   Group Key: t0.cat
   ->  Hash Join  (cost=3185.00..3197517.82 rows=98599882 width=12)
         Hash Cond: (t0.aid = t1.aid)
         ->  Seq Scan on t0  (cost=0.00..1833334.00 rows=100000000 width=8)
         ->  Hash  (cost=1935.00..1935.00 rows=100000 width=12)
               ->  Seq Scan on t1  (cost=0.00..1935.00 rows=100000 width=12)
(7 rows)

postgres=# SELECT cat,count(*),avg(ax) FROM t0 NATURAL JOIN t1 GROUP BY cat;
 cat |  count  |       avg
-----+---------+------------------
 nnn | 3845736 | 49.9122204644368
 ccc | 3842975 |  49.925316114681
 ddd | 3841209 |  49.934698930695
 aaa | 3848221 | 49.9265219996321
 kkk | 3843481 | 49.9232348058579
 fff | 3845484 | 49.9434949581975
 iii | 3846743 |    49.9227199719
 jjj | 3846076 | 49.9292863471066
 qqq | 3845646 | 49.9452133657041
 ttt | 3846725 | 49.9232478784215
 hhh | 3842519 |  49.926669332217
 zzz | 3847116 |  49.932075172461
 ooo | 3847927 | 49.9346102129966
 lll | 3848447 | 49.9278659066662
 www | 3846691 | 49.9537192102029
 bbb | 3845315 | 49.9281298849615
 ppp | 3850842 | 49.9149069792369
 eee | 3846285 |  49.931458202444
 xxx | 3845570 | 49.9577920754269
 ggg | 3845044 | 49.9409383291352
 uuu | 3845813 | 49.9295202543602
 rrr | 3847816 | 49.9341189910555
 vvv | 3849157 | 49.9253944163048
 yyy | 3843414 | 49.9364087656454
 mmm | 3848758 | 49.9033622681445
 sss | 3846990 | 49.9213589517182
(26 rows)

Time: 51031.084 ms

p2.* インスタンスの登場でPG-Stromを使用するための環境の準備はずいぶんお手軽になった。デプロイ即利用できるようAMIイメージも用意してみたので、まだPG-Stromを触った事がないという方は、ぜひ試してみて頂ければと思う。

*1:と言っても、Tesla K80なんですが…。

*2:予定なしとか悲しい事は言わないでほしいけど

2017年の開発ロードマップについて考える

あけましておめでとうございました。（やや出遅れ感）

新年という事で、この一年、どういった技術開発に取り組んでいきたいかをざーっと書き出してみる事にする。
これらのうち、いくつかはPostgreSQL本体機能の強化を伴うものであったりするので、ある程度計画的にモノゴトを進めないといけないワケで…。

PG-Strom v2.0

先ず最優先で取り組むのが、PostgreSQL v9.6への対応。
CPUパラレル実行と、新しいオプティマイザへの対応でかなり大きなアーキテクチャ上の変更を伴ったものの、全体としてはよりシンプルな設計に落とし込む事ができている。

ちなみに、現状だとこの程度までは動くようになっている。
集約演算がGroupAggregateとGpuPreAggの二段階に分解されており、GpuPreAggはGatherの配下で並列に動作している事に注目。

postgres=# EXPLAIN (ANALYZE, VERBOSE)
           SELECT cat, count(*), avg(aid), max(bid)
             FROM t0
            WHERE aid < 50000 and cid > 50000
            GROUP BY cat;

                                           QUERY PLAN
------------------------------------------------------------------------------------------------
 GroupAggregate  (cost=91050.50..91070.84 rows=26 width=48)
                 (actual time=1754.432..1755.056 rows=26 loops=1)
   Output: cat, pgstrom.sum((pgstrom.nrows())),
                pgstrom.favg((pgstrom.pavg((pgstrom.nrows((aid IS NOT NULL))),
                             (pgstrom.psum((aid)::bigint))))),
                max((pgstrom.pmax(bid)))
   Group Key: t0.cat
   ->  Sort  (cost=91050.50..91052.32 rows=728 width=48)
             (actual time=1754.381..1754.524 rows=910 loops=1)
         Output: cat, (pgstrom.nrows()), (pgstrom.pavg((pgstrom.nrows((aid IS NOT NULL))),
                                                       (pgstrom.psum((aid)::bigint)))),
                 (pgstrom.pmax(bid))
         Sort Key: t0.cat
         Sort Method: quicksort  Memory: 160kB
         ->  Gather  (cost=90938.54..91015.89 rows=728 width=48)
                     (actual time=1749.313..1753.732 rows=910 loops=1)
               Output: cat, (pgstrom.nrows()), (pgstrom.pavg((pgstrom.nrows((aid IS NOT NULL))),
                                                             (pgstrom.psum((aid)::bigint)))),
                       (pgstrom.pmax(bid))
               Workers Planned: 4
               Workers Launched: 4
               ->  Parallel Custom Scan (GpuPreAgg) on public.t0  (cost=89938.54..89943.09 rows=182 width=48)
                                                            (actual time=1670.139..1673.636 rows=182 loops=5)
                     Output: cat, (pgstrom.nrows()),
                             pgstrom.pavg((pgstrom.nrows((aid IS NOT NULL))),
                                          (pgstrom.psum((aid)::bigint))),
                             (pgstrom.pmax(bid))
                     Reduction: Local
                     GPU Projection: t0.cat, pgstrom.nrows(), pgstrom.nrows((t0.aid IS NOT NULL)),
                                     pgstrom.psum((t0.aid)::bigint), pgstrom.pmax(t0.bid), t0.aid, t0.cid
                     Outer Scan: public.t0  (cost=4000.00..87793.21 rows=2496387 width=12)
                                            (actualtime=14.663..274.534 rows=500187 loops=5)
                     Outer Scan Filter: ((t0.aid < 50000) AND (t0.cid > 50000))
                     Rows Removed by Outer Scan Filter: 1499813
                     Extra: slot-format
                     Worker 0: actual time=1724.195..1728.058 rows=182 loops=1
                     Worker 1: actual time=1570.952..1573.837 rows=182 loops=1
                     Worker 2: actual time=1738.205..1742.053 rows=182 loops=1
                     Worker 3: actual time=1569.055..1571.961 rows=182 loops=1
 Planning time: 0.907 ms
 Execution time: 1759.557 ms
(25 rows)

また、PG-Strom v2.0では、PostgreSQL v9.6へのキャッチアップだけではなく、いくつか目玉となる機能を準備中である。
一つは、これまで何度か紹介している SSD-to-GPU P2P DMA 機構。そしてもう一つは、BRINインデックスへの対応である。

SSD-to-GPU P2P DMA

SSD-to-GPU P2P DMA (NVMe-Strom) は、NVIDIA社製GPUのGPUDirect RDMA機構を利用したもので、PostgreSQLのデータブロックが格納されているNVMe-SSDのデータブロックからGPUへとダイレクトにデータ転送を行う。ファイルシステムを介する事によるオーバーヘッドや、RAMへの無駄なコピーが発生しないため、スループットを稼げるという特長がある。
現状では、GpuScanワークロード下においてNVMe-SSD 1個から成る区画からのデータ転送に対応しており、シングルプロセス性能で1.4GB/sのスキャン性能を出している。
PostgreSQL v9.6対応の過程で、GpuJoinやGpuPreAggの直下にテーブルスキャンが入る場合、これらのロジックはGpuScanがテーブルをスキャンするための関数を直接呼ぶように改良されているので、特別な事は何もしなくても『ストレージからデータを読んだ時点で既にJOIN/GROUP BYが完了している』という状態を作り出す事はできるはず。

PG-Strom v2.0に向けた課題はSoft-RAID0/1への対応。Linuxの場合、基本的には128KB単位で順番にストライピングがかかっているだけなので、技術的にはそう難しい話ではないと考えている。
DC用途向けに、PCIe x8スロット接続で5～6GB/s程度のSeqRead性能を持つNVMe-SSD製品が各社から出てきているので、計算上は、SSD二枚から全力でGPUにデータを流し込む事ができれば、GPUの持つPCIe x16スロットの帯域を飽和させられる事になる。

BRINインデックス対応

BRINインデックス自体はPostgreSQL v9.5から搭載されている機能で、特に時系列データのように

  1. ある一定範囲の値を持つデータが物理的な近傍に集まっている
  2. データの更新頻度が小さい
  3. データサイズが大きい

といった特徴を備えたデータセットに向いており、例えば、センサデータをPostgreSQLに収集して解析するといったワークロードに有効な機能。

BRINインデックス自体は、永安さんのこちらの記事が詳しいです。
pgsqldeepdive.blogspot.jp

PG-Stromとしては、搭載RAMが比較的小さなGPUを使うという事もあり、B-treeのようなランダムメモリアクセスを前提としたインデックスへの対応は厳しい。
ただ、条件句の評価はGPUの数千コアを使って並列処理が可能であるものの、インデックスの選択率が高くなると分の悪い勝負なので『このブロックは明らかに該当行なし』という事が分かっているなら、それを読み飛ばしたい。

GpuScanがBRINインデックスを理解し、必要のないブロックを読み飛ばす事ができるようになれば、例えばIoTのキーワードに絡めてセンサデータの集積・解析用途に、という使い方もできるハズである。
特に、PG-Stromはカラムナーを前提としたDWHではないので、生データをそのまま処理させても高速化できるという点は強みになるだろう。

PL/CUDA

PL/CUDAに関しては、言語バインディングに関してする事は（できる事は）多くないので、その周辺領域を拡充していきたい。

一つは、PostgreSQLにおける可変長データの1GB越え。
現状、全ての可変長データの基盤となっている varlena 構造は、最大でも1GBまでのデータしか持てないため、PL/CUDA関数の引数としてarray-matrixを渡す時には、例えば問題領域をうまく分割してデータサイズを1GB未満に抑えてやらないといけない。

昨年10月のCBI学会で発表した研究でも、1000万件の化合物データ（1.3GB）をロードするにはサイズが大きすぎたので、安全マージンも見て4分割した上でGPU側へロードしている。

しかし、昨今のGPUでは10GB近く、あるいはそれ以上のメモリを搭載するのが常識的になりつつあり、問題領域を1GB以内に抑えねばならない、、、というのはユーザにいかにも不都合である。

先にpgsql-hackersにデザインプロポーザルを投げたところ、Robert Haasから『varlenaとは別の体系で可変長データを保持するフォーマットを作成すべき』とサジェスチョンがあり、自分もこの方針には同意。3月のcommit-festまでにはパッチを投稿し、2018年リリース予定のPostgreSQL v11でのマージを目指したい。

もう一つは、現状、複雑な計算ロジックを個々のユーザ毎に書かねばならないという点である。
2017年1月時点でPL/CUDAを実証できたワークロードとしては、k-NN法類似度サーチや、k-meansクラスタリングがあるが、例えば MADLib のような統計解析パッケージで提供されているアルゴリズムの、全部とは言わないまでも、使用頻度が高く計算負荷の高いものをGPUで計算するようパッケージ化できれば、ユーザの裾野はより広がるだろうし、仮にカスタマイズが必要となっても骨格となるアルゴリズムのGPU実装が既に存在する事で省力化が可能となるハズである。

PG-Strom v3.0へ向けた種まき

In-memory Columnar Cache

NVMe-SSDとの密連携の他にI/O系処理を高速化する方策として、通常のPostgreSQLテーブル（行形式）の脇に、予めDMAバッファ上に列形式にデータを再編したキャッシュを持たせる機構を考えている。
このキャッシュはBackground-workerにより非同期で作成され、そのため、スキャンする区間のうち一部領域しか列形式のキャッシュが構築されていないかもしれない。しかし、その場合でもPG-StromはGPUで行形式データを扱えるので（多少のパフォーマンス差に目をつぶれば）大きな問題とはならない。

データ本体とキャッシュを別に持つ場合、必ず一貫性制御が複雑で頭の痛い問題として立ちはだかる。
アイデアとしては、これを避けるために ALL_VISIBLE フラグが 1 であるブロックのみを列形式キャッシュに持たせる。
ALL_VISIBLE=1であるブロックは、MVCC制御に関わらず、全てのレコードが全てのトランザクションから可視である事が保証されている。そのため、複雑な同時実行制御に頭を悩ませる必要はなく、全てのレコードの内容を単純に列データとして展開すればよい。

問題は、PostgreSQL側でテーブルが更新され、ALL_VISIBLE フラグが 0 にクリアされた時のinvalidation処理である。
現状、ここにフックを挟む事はできないので、PostgreSQL側の機能強化を行う必要がある。
デザインプロポーザルを出し、この処理を行うにふさわしい場所とフックの仕様を固めていきたい。

GpuSort(+LIMIT)

PG-Strom v2.0では、実はGpuSort機能の廃止を予定している。これは、ソートという問題の性質上、ある程度問題規模が大きくならないとGPUによる処理時間メリットが出てこない一方で、一度にGPUでソートを行う件数が多くなればなるほど、初期データのローディングに時間がかかるようになり、非同期・多重処理のメリットを得にくいからである。
そのため、GPUでソートを行うセグメントサイズを小さくして単位ローディング時間を短くする一方で、CPUでのMergeSort処理の割合が大きくなるか、それとも、セグメントサイズを大きくするかというジレンマに悩まされてきた。
（で、最終的にはそれほど速くならない事が多かったり・・・。）

根本的な原因は、GPUで処理を行ってもソートではデータ件数を減らす事が不可能な点にある。なので、GpuSortを有効に活かすには、何がしか『データ件数を減らせる』パターンでのソートに限った方が利口だ。
例えば、LIMIT句で『上位xx件を取り出す』という事が明らかな場合に限り、GpuSortを使用するというパターンであれば十分に効果を発揮する事ができるだろう。
これにはPostgreSQL本体側で、『LIMIT句でxx件のデータが要求されますよ』という事を下位のノードに伝えてやる仕組みが必要だが、PostgreSQL v10 向けにパッチを出しており、ある程度レビューも進んでいるので間に合うだろうとは踏んでいる。

こんな感じで2017年の開発ロードマップについて考えてみたが、さてさて、大晦日に振り返ってどの程度きちんとやれているでしょうか。