GTCで喋りました

という訳で、GTCで喋ってきました。

www.slideshare.net

今までの発表とは少し趣を変えて、PG-Stromそのものの説明よりも、現実世界のワークロードを実行するときにどういった使い方があり得るか、どういった効果が得られるかに主眼を置いた内容。

f:id:kaigai:20160408081720p:plain
伝統的にDBMSは、正に字面の通り、データベースをマネジメントするソフトウェアである訳なので、インデックスを張って目的のレコードを高速に取り出したり、テーブル同士を結合する事は得意だが、(一応SQLの構文で書く事はできるとはいえ)大量の計算をこなすには向いていない。
少なくとも、データ解析専用に設計されたツール(例えばR言語)を使って計算させた方が高速で、しかも統計解析系のパッケージも揃っているので、普通はデータベースからデータをエクスポートした上でデータ解析処理を行う。

では、PG-Stromを入れることで計算処理が早くなれば、この前提が変わるか?というのが今回のお題。

■ DB内で解析処理を行う事のメリット

  • サーバ~クライアント間でデータを移動せずにすむ。数GB単位の大きさになると無視できない。
  • 通常はクライアント側よりも強力なサーバハードウェアの計算能力を使用できる。
  • 常に最新のデータセットに対して解析処理を行うことができる。

■ DB内で解析処理を行うことのデメリット

  • サーバ側が十分な計算能力を持っている必要がある。
  • DBMSがデータ解析処理に対応するよう設計されている必要がある。
  • 専用APと比べると、ライブラリ類の整備が不十分。

ライブラリ類の充実はともかくとして、計算能力に関してはいかほどのものか。

f:id:kaigai:20160408081727p:plain

今回の実験のターゲットは、創薬に関連したワークロードの一つで、大量の化合物の中からできるだけ特性の異なった化合物を選択するというもの。これにより、実際に化合物を合成して特性を調べるなど、"お金のかかる" ステップを少なくしたいというモチベーションがある。

f:id:kaigai:20160408081740p:plain

使用したアルゴリズムはMAX-MIN法と呼ばれるもの。実際には何個かのクエリを繰り返し実行するのだが、ワークロードの中核、最も計算が重い部分は赤点線で囲った距離計算の部分。
化学物質の特性を数値パラメータ化し、比較対象との"距離"を計算する。
基本的にはこの距離が大きな化学物質を順に選択することで、ある一群からピックアップした k 個の化学物質の特性ができるだけ他と異なるものになる。

f:id:kaigai:20160408081748p:plain
これを、SQLによるデータエクスポート+R言語でのローカル計算、SQLで実行+R言語で結果参照 with/without PG-Strom の3パターンでそれぞれトライしてみた。
結果は、PG-Stromなしのパターンが圧倒的に遅く、R言語のローカル実行と、SQL+PG-Stromのリモート実行が概ね同じという結果になった。

そのため、現状でも『わざわざデータをEXPORTしてローカルで計算しなくても、SQLで全部計算させて結果だけ取ってきた方がデータ管理が楽ですよ』というのは言えるだろう。


一方で、他の並行セッションでは、よりHPCに近く計算能力を必要とする人工知能の研究成果が発表されており、『SQLを高速に処理する』がためにGPUを使用するPG-Stromの設計では専用にチューンしたアプリケーションに太刀打ちすることはできない。
(する必要があるのかどうか、という議論は当然あるが)

しかし、データに最も近い場所でCUDAプログラムを実行しうるプラットフォーム機能を既に有しているというアドバンテージを活かすのであれば、OLTP/OLAPという伝統的DBの世界から外れたところのワークロードについて考えてみる価値はあるだろう。おそらく、必要な要素技術は既に揃っている。

f:id:kaigai:20160408085151p:plain

検討してみたいのは、前回のエントリでも書いたアイデアで、CUDAのロジックをSQLで記述するための機能。

PostgreSQLでは、SQL関数を記述する言語を拡張モジュールによって定義する事ができる。
したがって、CUDAによって記述されたSQL関数を実現するのはそれほど難しい話ではない。

SQL上で hogehoge(matrix, vector) などと定義された関数であれば、実行時には、第一引数に行列をコピーしたバッファのポインタを、第二引数にベクトルをコピーしたバッファのポインタを渡すようにすればよい。フラットな単純配列に対してCUDAカーネルを起動すれば、あとは普通のHPC屋さんがやっている内容になる。

現状のSQL->CUDAへのコード変換の場合、SQLに由来するNULLチェックや四則演算の都度発生するオーバーフローチェックが入っているため、計算処理に重きを置くのであれば相当に効率の悪いコードになっている。

何から何までCUDAのコードを書かせるのはアレだが、R言語などでもパッケージ化されていて比較的良く使われる関数、例えば kmeans() 相当がCUDAのフルスピードで実行できるようになったらどうだろうか?

SQLでkmeans法のアルゴリズムを記述した場合、繰り返し処理の記述性が高くなく無駄な処理が入っていることもあり、現状では PG-Strom を用いてもRのkmeans()関数の1.5倍程度の時間を要している。(PG-Strom無しだと実行完了を諦めるレベルだが)

CUDAのフルスピードで行列演算を実行できれば、計算部分の処理速度を現状から更に一桁向上させる程度はできるだろう。
RDBMS的にはニッチもニッチでいい所の機能だが、『明らかに高速化の効果を実感できる』機能として、前倒しトライしてみるのも良いかなと思えてきた。

GTCに来ています

今年もサンノゼで開催されている GPU Technology Conference 2016 に参加しています。

# なお、当方の発表『In-Place Computing on PostgreSQL: SQL as a Shortcut of GPGPU』は木曜日の予定

f:id:kaigai:20160405081234j:plain
キーノートでは、NVIDIA社CEOのJen-Hsum Huang氏より"ディープラーニング向けプロセッサ"と銘打って、新アーキテクチャPascalを搭載したTesla P100が発表に。また、CUDA8.0のリリースが6月である事も発表された。
(一方で、事前に噂されていたPascalベースのコンシューマ向けGTXモデルへの言及はなかった)

f:id:kaigai:20160405102340j:plain

このPascalアーキテクチャベースのTesla P100であるが、Inside Pascal: NVIDIA’s Newest Computing Platformに記載のある通り、

...等々のお化けプロセッサである。

プログラムを書く上で、Maxwell世代から変わっている点、新機能が追加になった点がある。

  • SM(Streaming Multiprocessor)あたりのCUDA Core数が減っている

128個 -> 64個になった。一方でSM数は56個に増えており、これでトータル3584CUDA Coreという訳だが、SMあたりのレジスタファイル容量(256KB)は据え置き、共有メモリ容量(64KB)は33%増加*1であるので、システム全体として見ると、大幅な増強となっている。

また、SMあたりのコア数が少なくなった一方で、レジスタ、共有メモリの容量が据え置きになっているという事は、ある特定のSMに対して一度に投入できるタスクの数が増え、それにより、例えばDRAMアクセス待ちの間に他のタスクにスケジュールして計算コアの使用率を上げる事ができると思われる。

PG-Stromのようにnone-coalescedなメモリアクセスが多いワークロードでは嬉しいかも。

PG-Stromには嬉しい新機能。現状、PostgreSQL数値計算系の関数は全て64bit浮動小数点で、標準偏差や分散、共分散などを集計する時には、内部的には cmpxchg 命令を使った atomic 演算のエミュレーションを行わざるを得ないので、H/Wによりネイティブの atomic 演算がサポートされれば集計系の性能向上が期待できる。

  • プリエンプションの対応

仮想化環境での利用を見込んだ機能か?上述のSMに投入できるブロック数が増えたという事も併せて考えると、Pascal世代ではGPUの同時並行利用が大いに促進されるような気がする。

  • デマンドページング

従来は、予め必要となるバッファサイズを予測してCUDA Kernelを起動する前にデバイスメモリのアロケーションが必要であったが、これにはある程度マージンを取る必要があるので、実際には無駄になる領域が少なからずあった。GPU側でもPage Fault→デバイスメモリ割当てという流れが可能になる事で、リソースの有効利用が可能となる。
一方で、Page Faultのコストは無視できないほど高いので、cudaMemPreFetchAsync()などCUDA8.0で対応となる新APIにより、プログラムがヒントを与えてやる必要がある。

  • HBM2

一目で分かるように、スループットが大幅に引き上げ。現行世代の3倍弱。
ただし、これはメモリアクセスがcoalescedな場合の性能向上で、none-coalescedなメモリアクセスはレイテンシが律速要件になるため、現行世代と比較して10%程度の向上との事。

行指向データ構造を使うPG-Stromとしてはちょっと悲しい所だが、対策としては、例えばDRAMからデータを一旦共有メモリにロード(ここはcoalescedなアクセスで)。その後の共有メモリへのアクセスはL1キャッシュと同じレイテンシなので、そんなに気にはならないはず。

SMあたりのCUDAコア数が減った事で、こういう最適化も考えられる。


さて、会場でSさんとHさんに『PG-Stromの上でMachine Learningのアプリとか動かせるようにならないんですかね~?』と無茶振りをされる。

現状、"PostgreSQLと同じ使用感"で使える事を優先するために、本来は効率の悪い行指向データをリッチな計算リソースで無理やり並列処理しているので、専用のアプリケーションでの処理と比べるとどうしても分が悪い。

ただ『データのある場所で計算を行う』というのは正しい方向性なので、SQLで数式を書いたら自動的にコード生成⇒GPU実行、という流れに拘らなければ、SQLクエリ実行の過程でデータ密度の高いデータ構造を作り出し、これをGPUで実行するMachine Learning用の(でも、何か他の数値計算の)アプリケーションに渡してやるという事はそんなに不自然ではないのかな、とは思う。

例えば、32bit浮動小数点型の2次元配列・NULLなしというデータ構造であればぎっしりデータが詰まっており、先頭からのオフセットだけでアクセスすべきデータを特定できる。つまり、行指向ならではの不都合はない。これを仮に matrix 型と名付けて、

SELECT my_deep_leraning(SELECT make_matrix(x0, x1, x2, ..., xN)
                          FROM data_source
                         WHERE 抽出条件
                         GROUP BY グループ化条件);

みたいなクエリで実行できれば、、、、って事にはなるだろう。

ただ、当然に考えねばならん課題もあり、
・生成するMATRIXがGPUのRAMサイズに収まらない場合
・粗行列の場合、本当に単純な配列表現でいいのか?
・一行に対して複数のCUDA Kernelで処理する事になるので、PG-Stromにとってはwhat's new

面白いアイデアなので、もう少し考えてみたいところではある。

*1:たぶんブロックあたりのサイズの事を言っている。SMあたり容量ならMaxwellの上の方は96KBなので

SSD-to-GPU Peer-to-Peer DMAとバッファ管理(その2)

前回の続き。

PCI-E接続のSSDからのP2P DMAでCPU/RAMを介さずにGPU RAMへデータを転送するという要素技術自体は目新しいものではない。

かつてFusion-io(現: SunDisk)もやっていたし、NVMe規格に準拠したものであれば標準のドライバに少し手を加えてP2P DMAを実行する事も可能である。実際にやっている人もいる。

問題は、PostgreSQLでの利用シーンを考えた時に、常にストレージからのP2P DMAを実行すればよい、というワケではなく、バッファにキャッシュされたデータの扱いや並行プロセスとの相互作用を考慮して、矛盾なく、かつ、性能的にもリーズナブルである方法を考えねばならない。

先ず、大前提として考えねばならないのが、PostgreSQLの追記型アーキテクチャとMVCC(Multi-version concurrency control)の仕組み。
かいつまんで言うと、各レコードには『そのレコードを作成した人のトランザクションID(xmin)』と『そのレコードを削除した人のトランザクションID(xmax)』という値が記録されており、参照側トランザクションIDとの大小関係や、これらのトランザクションがcommitされているか否かといった情報を元に、レコードの可視/不可視を判定する。

これらのルールはトランザクション分離レベル(READ COMMITTEDやREPEATABLE READ)によって異なるが、それなりに複雑なロジックを処理するのでできれば避けたい。(さらに言えば、同時並行で動いているトランザクション状態にも依って判定が変わるため、GPUで処理する事はできない。)

そこで使われるのが visibility map と、all_visible フラグ。

f:id:kaigai:20160220123541p:plain

PostgreSQLでは8KB~32KBの大きさのブロック単位でデータを管理しており、ブロックに包含されるレコードが全て、あらゆるトランザクションから明らかに可視である時には all_visible フラグがセットされる。このフラグが立っている時には、個別レコードのMVCCチェックは必要ない。
さらに、これはブロック自体とは別管理の visibility map(VM) と呼ばれるデータ構造と連動しており、VMのビットが立っているブロックは、(ブロック自体を読み込まなくても)all_visibleフラグがセットされている事が分かる。*1

そもそもGPUではall_visible=0であるブロックからレコードを抽出できないので、all_visible=0であるブロックはCPU側でレコードを抽出してRAM->GPUへと転送しなければならない。したがって、この時点でSSD-to-GPU P2P DMAの対象にできるのはall_visible=1のブロックである事が分かる。*2

VMの状態から次に読もうとするブロックがGPUでの処理が可能である事が分かった。次に、このブロックが既にPostgreSQLのバッファに載っているかどうかを確認する。
NVMe SSDは確かに高速なデバイスではあるが、RAMと比較すると文字通り"桁違いの"遅さであるので、同じPCI-Eバスを介してデータ転送を行うのであれば、普通にRAM->SSDへのDMAを行った方が合理的である。

したがって、SSD-to-GPU P2P DMAの対象となるブロックは ①all_visibleフラグが立っている ②まだshared_bufferにロードされていない という条件を満たすものである事が分かる。

f:id:kaigai:20160220123545p:plain

さて、並行するプロセスが介在する場合を考える。ある時点で対象となるブロックがshared_bufferに載っていない事が分かり、P2P DMAのリクエストをキューに入れたとする。しかし、それとはお構いなしに、並行プロセスが当該ブロックをOSからロードし、更新処理を行ってしまうかもしれない。
ただし、この事自体は問題とはならない。更新されたレコードは、P2P DMAを行っているトランザクションが開始した時点では未だコミットされておらず、あたかも存在しないものであるかのように扱っても構わないからである。
したがって、OS buffer/Storageに保持されている古いバージョンのブロックをGPUへ転送し、そこで処理を行っても問題とはならない。

f:id:kaigai:20160220123555p:plain

注意しなければならないのは、並行プロセスによって更新されたレコードが、P2P DMAの完了までにOSに書き戻される場合。

f:id:kaigai:20160220144953p:plain

PostgreSQLのストレージ書き込み(smgrwrite)はBuffered Writeなので、書き出された内容はまずOSのPage Cacheに保持される。これは単なるメモリコピーであるが、Page CacheからGPU RAMへとDMA転送を行っている途中にPage Cacheが更新され、旧バージョンと新バージョンが混じり合った状態のデータがGPUに転送されてしまうと最悪である。

ファイルシステムにもよるが、write(2)システムコールの処理中は inode->i_mutex により排他ロックを取る事になっているようなので、不完全な状態のデータを使用する事を避けるには、P2P DMAの実行開始*3から転送完了までの間は、NVMeドライバの側でも inode->i_mutex を取ってやらないとダメ。

さらに、inode->i_mutex で排他処理を行った場合でも、P2P DMAの実行より前に、並行プロセスによって更新済みページが既にOSバッファに書き出されたという場合が考えられる。このようなケースでは、GPU側でまずブロックの all_visible フラグをチェックし、GPU RAMに転送された時点でこれが all_visible=0 にクリアされていた場合にはいったん諦め、CPU側でMVCCチェックと改めてRAM上に確保したDMAバッファにレコードを積んでおく事を要求する。

そもそも、P2P DMAリクエストから転送完了までのわずかな時間の間に、並行プロセスが当該ブロックをOSから読み出し、更新を行い、それをまたOSバッファに書き戻すという一連の処理を実行する頻度はかなり低いと想定される上に、更新済みブロックは並行プロセスの手で既に shared_buffer にロードされているので、この再実行によって新たにI/Oの負荷が発生する事はない。

気を付けなければならないのは、OSのPage Cacheへの書き出しとDMAの処理がカチ合わないよう、ここだけはきちんと排他処理を踏んでやる事である。

*1:これはIndexOnlyScanでも使われており、PG-Stromとは無関係に元々ある機能

*2:とはいえ、一般的な業務アプリでは90%の行は更新されない(by SAP)らしいので、通常はall_visible=1となるブロックが大半であろう。

*3:リクエストキューに積んだ時点である必要はない

SSD-to-GPU Peer-to-Peer DMAとバッファ管理(その1)

昨年の暮れ、JPUGカンファレンスのLTで『SQL+GPU+SSD=∞』と題したスピーチを行った。

www.slideshare.net

これはかいつまんで言えば、ストレージからデータをCPU+RAMへとロードするより前に一旦GPUへとデータを転送し、そこで不要なデータを削ぎ落してからCPU+RAMへと渡してやる事で、CPU負荷の軽減とRAMの有効活用が計れるというアイデアである。
実装としては、PCI-Eバイス間でのP2P DMA機能を利用する事によってNVMe SSDの特定ブロックからGPU RAM上の特定の領域へDMAを実行するというものなので、ここは別に新しくも何ともない。

以下の図は、従来の仕組みにおけるデータの流れを示したもの。
f:id:kaigai:20160214001237p:plain
SSDから読み出されたデータは先ずCPU+RAMにバッファされ、それをGPUに送出してJOIN処理を実行、結果を受け取るという一連の流れである。しかし、①SSD=>RAMにデータが転送される時、本来は検索条件に合致せず削除されるはずのタプルや、問合せの中で参照されないカラムも含まれている。つまり、ゴミデータもRAM上のバッファを消費し、もしかしたらその煽りで別のキャッシュされていたブロックが弾き飛ばされてしまう可能性だってある。

もしPG-StromがSSD-to-GPU P2P DMA機能に対応すればどのように変わるか?
f:id:kaigai:20160214001242p:plain
SSDから読み出されたデータは、CPU+RAMへの転送より前に、先ずGPUへと転送される。ここで検索条件や参照列のチェックを行い、先ず、不要な行と列は削除する。場合によっては、SSDから直接読み出したテーブルと結合する事になるInner側のテーブルが既にCPU側から送られており、不要行・列の削除に留まらず、JOINまで行った状態でCPU+RAMへ結果を転送できるかもしれない。
こうすると、CPU+RAMの視点から見た時、『SSDからデータブロックを読みだしたと思ったら、なんとJOINまで既に終わっていたでござる』という状況を作る事ができるようになる。

ただ、ここで考えねばならないのがバッファ管理。
PostgreSQLの全てのブロックの内容がストレージ(NVMe SSD)と常に同期されていれば話は簡単なのだが、そうは問屋が下ろさない。

PostgreSQLはデータブロックの書き出しをbuffered writeで行うので、ストレージの内容よりも新しいデータが、PostgreSQL自身のバッファ(shared buffer)、あるいはOSのバッファ(page cache)に存在する可能性がある。
したがって、バッファされていない、即ちストレージに最新のデータが存在する場合にはSSD-to-GPU P2P DMAを行うとしても、バッファされているデータとの整合を考慮しなければならない。
つまり、GPUを使おうとするプログラムは、バッファリング状態を確認した上で適切なDMA元を指定してやる必要がある。

f:id:kaigai:20160214001305p:plain

先ず、目的とするブロックがPostgreSQLのshared buffer配下に存在するか否かの確認。これは BufTableLookup() 関数を用いれば確認する事ができ、既に載っていれば、shared bufferをソースとするDMAを実行する。
既にCPU+RAM側で保持しているデータに対しては、GPUでデータの振るい落としを行おうがそもそもRAM使用量の節減効果はない事と、そもそもRAMはSSDよりも速いデバイスだから、という事である。

PostgreSQLのshared bufferに載っていない場合、Linuxのpage cache状態の確認はユーザプロセスからは行えないので、ドライバ側で適切にハンドリングする必要がある。これも find_get_page() 関数を用いれば確認する事ができ、既に載っていれば、page cacheをソースとするDMAを実行する。
この2パターンはどちらも本体RAMをソースとするDMAで、それに該当しない場合、NVMe SSDの該当ブロックをソースとするDMAを実行する事になる。

ちょっと辛いのが、『特定ファイルの特定オフセット位置』をブロックデバイス上のセクタ番号に変換するためのインターフェースがVFSで規定されていない事。なので、NVMe SSDドライバにSSD-to-GPU対応機能を追加するとしても、利用可能なファイルシステムを何個か絞った上で、ファイルオフセット⇒ブロック番号への変換関数を呼びだして対処する必要がある。*1
Ext4ならext4_get_block、XFSならxfs_get_blocks関数なんてものが定義されているので、これらを呼び出せばそう難しいものではないハズだが…。

もう一つ考慮しなければならないのが、競合する更新処理をどのように捌いていくのか。これはエントリを分けて記す事にしたい。

*1:LVMやSoft-RAIDを使用する場合はもう一段階の変換が必要か

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非対応

俺様スキャンの並列実行

PostgreSQL v9.6からはパラレルスキャンが導入される事になっている。

この機能をざっくり説明すると

  1. 共有メモリ上に『次に読むブロックの番号』という状態を作っておく。
  2. Gatherノードが複数のワーカープロセスを起動する。
  3. 各ワーカーで実行されるSeqScanが『次に読むブロックの番号』をアトミックにインクリメントしながらシーケンシャルスキャンを行う。
  4. 結果、個々のワーカーは本来読むべきテーブルの内容を一部しか返却しない。しかし、各ワーカーの実行結果を全てマージするとテーブルの内容は全て読み出されている。(しかも並列に)

という代物である。
読み出すブロック番号を共有メモリ上の情報から得ている以外は、全てシングルスレッド実行と同じコードを使用しているところがポイントである。

で、これと同じ事をForeignScanやCustomScanを使用して実装したい場合、並列対応版とそうでないものの差は、①この共有メモリ上に作られる『次に読むブロックの番号』(や、それに類する情報)を作るための仕組みと、②この情報に基づいて読み出す範囲を決定するという部分だけである。

②に関してはモジュール側で実装する部分の範疇なので特にインターフェースの拡張は必要ないが、①に関しては、共有メモリを獲得・初期化するタイミングでモジュール側を呼び出してやる必要がある。

その呼び出しポイントは、以下の関数。

  • ExecParallelEstimate ... 共有メモリセグメントの長さを決定する
  • ExecParallelInitializeDSM ... 共有メモリセグメントの初期値を設定する
  • ExecParallelInitializeWorker ... ワーカー側で実行ノードの初期化を行う

そのそれぞれに対応するCustomScan/ForeignScanのコールバックを定義してみた。

CustomScanはこんな感じ。

Size (*EstimateDSMCustomScan) (CustomScanState *node,
                               ParallelContext *pcxt);
 → 必要な共有メモリのサイズを返す

void (*InitializeDSMCustomScan) (CustomScanState *node,
                                 ParallelContext *pcxt,
                                 void *coordinate);
 → coordinateに割り当てられた共有メモリを初期化

void (*InitializeWorkerCustomScan) (CustomScanState *node,
                                    shm_toc *toc,
                                    void *coordinate);
 → worker側で共有メモリの内容に基づいて、実行ノードの
    (追加的な)初期化を行う。

ForeignScanもこんな感じ。全く同様

Size EstimateDSMForeignScan(ForeignScanState *node,
                            ParallelContext *pcxt);
void InitializeDSMForeignScan(ForeignScanState *node,
                              ParallelContext *pcxt,
                              void *coordinate);
void InitializeWorkerForeignScan(ForeignScanState *node,
                                 shm_toc *toc,
                                 void *coordinate);

このコールバックを使って、file_fdwを並列実行対応にしてみた。
diff差分を見てみると、修正規模は300行強。非常に小さな手間で並列化ができた。

postgres=# set max_parallel_degree = 3;
SET
postgres=# explain analyze select * from test_csv where id % 20 = 6;
                                  QUERY PLAN
--------------------------------------------------------------------------------
 Gather  (cost=1000.00..194108.60 rows=94056 width=52)
         (actual time=0.570..19268.010 rows=2000000 loops=1)
   Number of Workers: 3
   ->  Parallel Foreign Scan on test_csv  (cost=0.00..183703.00 rows=94056 width=52)
                                  (actual time=0.180..12744.655 rows=500000 loops=4)
         Filter: ((id % 20) = 6)
         Rows Removed by Filter: 9500000
         Foreign File: /tmp/testdata.csv
         Foreign File Size: 1504892535
 Planning time: 0.147 ms
 Execution time: 19330.201 ms
(9 rows)

topで見ても、worker(ForeignScan)×3とmaster(Gather)×1がぶん回っているのが分かる。
f:id:kaigai:20160129011848p:plain

TargetListの計算をGPUで行う。

正月休みの宿題だった機能を実装できた。(注:ちゃんと動くとは言っていない)

PG-Stromを使って数式の評価をGPUにオフロードする場合、WHERE句やJOIN..ON句のオフロードには対応していたものの、TargetListに複雑な演算式を含む場合、これは完全にCPU側で処理せざるを得なかった。
これはPostgreSQLオプティマイザがこれらの数式をPlanノードにアタッチするタイミングの問題で、別に本質的な問題がある訳ではない。(実際、コミュニティの側でも改善に向けた動きはある)

例えば、以下のようなクエリは数式の評価をGPUで実行できた。

-- 任意の二点間の4次元空間上の距離が 10 未満であるペアを出力
SELECT a.id a_id, b.id b_id
  FROM test a, test b
 WHERE a.id > b.id AND sqrt((a.x1 - b.x1)^2 +
                            (a.x2 - b.x2)^2 +
                            (a.x3 - b.x3)^2 +
                            (a.x4 - b.x4)^2) < 10;

EXPLAIN文で確認すると、以下のように、GpuNestLoopのJoinQualに上記のWHERE句が収まっていることが分かる。

# EXPLAIN
  SELECT a.id a_id, b.id b_id
    FROM test a, test b
   WHERE a.id > b.id AND sqrt((a.x1 - b.x1)^2 +
                              (a.x2 - b.x2)^2 +
                              (a.x3 - b.x3)^2 +
                              (a.x4 - b.x4)^2) < 10;

                        QUERY PLAN

------------------------------------------------------------
 Custom Scan (GpuJoin)  (cost=14050655.37..25161766.48 rows=1111111111 width=8)
   Bulkload: On (density: 100.00%)
   Depth 1: GpuNestLoop, JoinQual: ((id > id) AND (sqrt(((((((x1 - x1))::double precision ^ '2'::double precision) + (((x2 - x2))::double precision ^ '2'::double precision)) + (((x3 - x3))::double precision ^ '2'::double precision)) + (((x4 - x4))::double precision ^ '2'::double precision))) < '10'::double precision))
            Nrows (in/out: 1111111.13%), KDS-Heap (size: 8.11MB, nbatches: 1)
   ->  Custom Scan (BulkScan) on test a  (cost=0.00..1637.00 rows=100000 width=20)
   ->  Seq Scan on test b  (cost=0.00..1637.00 rows=100000 width=20)
(6 rows)

なお、テーブル test は 4次元空間上の特定の点を表現するとして、以下のように定義。

postgres=# \d test
     Table "public.test"
 Column |  Type   | Modifiers
--------+---------+-----------
 id     | integer | not null
 x1     | real    |
 x2     | real    |
 x3     | real    |
 x4     | real    |
Indexes:
    "test_pkey" PRIMARY KEY, btree (id)

しかし、同じ式を含んでいても、このパターンだと無理だった。

-- 任意の二点間の4次元空間上の距離を算出してIDのペアと共に出力
SELECT a.id a_id, b.id b_id, sqrt((a.x1 - b.x1)^2 +
                                  (a.x2 - b.x2)^2 +
                                  (a.x3 - b.x3)^2 +
                                  (a.x4 - b.x4)^2) dist
  FROM test a, test b
 WHERE a.id > b.id;

同様にEXPLAIN文で確認すると、a.id > b.id の部分しかGPUで実行されていない。計算ロジックとして重いのは距離計算の部分であるにも関わらず、だ。

# EXPLAIN
  SELECT a.id a_id, b.id b_id, sqrt((a.x1 - b.x1)^2 +
                                    (a.x2 - b.x2)^2 +
                                    (a.x3 - b.x3)^2 +
                                    (a.x4 - b.x4)^2) dist
            FROM test a, test b
           WHERE a.id > b.id;
                                    QUERY PLAN
-----------------------------------------------------------------------------------
 Custom Scan (GpuJoin)  (cost=776405.37..167443072.02 rows=3333333333 width=40)
   Bulkload: On (density: 100.00%)
   Depth 1: GpuNestLoop, JoinQual: (id > id)
            Nrows (in/out: 3333333.20%), KDS-Heap (size: 8.11MB, nbatches: 1)
   ->  Custom Scan (BulkScan) on test a  (cost=0.00..1637.00 rows=100000 width=20)
   ->  Seq Scan on test b  (cost=0.00..1637.00 rows=100000 width=20)
(6 rows)

では、どのタイミングで計算が行われているのか?これはEXPLAIN VERBOSEで確認することができる。

# EXPLAIN VERBOSE
  SELECT a.id a_id, b.id b_id, sqrt((a.x1 - b.x1)^2 +
                                    (a.x2 - b.x2)^2 +
                                    (a.x3 - b.x3)^2 +
                                    (a.x4 - b.x4)^2) dist
    FROM test a, test b
   WHERE a.id > b.id;

                        QUERY PLAN

-------------------------------------------------------------------------
 Custom Scan (GpuJoin)  (cost=776405.37..167443072.02 rows=3333333333 width=40)
   Output: a.id, b.id, sqrt(((((((a.x1 - b.x1))::double precision ^ '2'::double precision) + (((a.x2 - b.x2))::double precision ^ '2'::double precision)) + (((a.x3 - b.x3))::double precision ^ '2'::double precision)) + (((a.x4 - b.x4))::double precision ^ '2'::double precision)))
   Pseudo Scan: a.id::integer, a.x1::real, a.x2::real, a.x3::real, a.x4::real, b.id::integer, b.x1::real, b.x2:
:real, b.x3::real, b.x4::real
   Bulkload: On (density: 100.00%)
   Depth 1: GpuNestLoop, JoinQual: (a.id > b.id)
            Nrows (in/out: 3333333.20%), KDS-Heap (size: 8.11MB, nbatches: 1)
   Features: format: tuple-slot, bulkload: supported
   Kernel Source: /opt/pgsql/base/pgsql_tmp/pgsql_tmp_strom_31576.1.gpu
   ->  Custom Scan (BulkScan) on public.test a  (cost=0.00..1637.00 rows=100000 width=20)
         Output: a.id, a.x1, a.x2, a.x3, a.x4
         Features: format: tuple-slot, bulkload: supported
   ->  Seq Scan on public.test b  (cost=0.00..1637.00 rows=100000 width=20)
         Output: b.id, b.x1, b.x2, b.x3, b.x4
(13 rows)

やや煩雑な出力結果となってしまっているが、見るべきポイントは GpuJoin の Output: と Pseudo Scan: の差分。
GpuJoinはGPUでのJOIN処理結果として、Pseudo Scan: の定義に従ってレコードを生成し、これをCPU側へ返却する。
一方、それを受け取ったGpuJoinノードは、より上位へ結果を返すために Output: の定義に従ってレコードの内容を書き換える。これをProjection処理と呼ぶ。
Pseudo Scan == Output であれば実際にはProjectionは必要ではないが、上記のように複雑な計算式を伴うケースであれば、ScanやJoin処理よりもProjectionが処理時間の支配項となってしまう。


では、どうすればよいか?
CPUでのProjectionが支配項となってしまうなら、Projectionが発生しないよう、予めGPU側で計算した結果を書き戻してやればよい。これはPostgreSQLオプティマイザがパスを検討する時点ではどういった数式を処理すべきか判断できないため、planner_hook を使って無理やりプラン木を書き換える事になる。(v9.7辺りでもう少しマシな手段が使えるようになればいいけどー)

で、正月休みにシコシコと作業をした結果が以下の通り。

# EXPLAIN VERBOSE
  SELECT a.id a_id, b.id b_id, sqrt((a.x1 - b.x1)^2 +
                                    (a.x2 - b.x2)^2 +
                                    (a.x3 - b.x3)^2 +
                                    (a.x4 - b.x4)^2) dist
    FROM test a, test b
   WHERE a.id > b.id;

                        QUERY PLAN

-------------------------------------------------------------------------
 Custom Scan (GpuJoin) on public.test a  (cost=777399.00..167444065.65 rows=3333333333 width=40)
   Output: a.id, b.id, (sqrt(((((((a.x1 - b.x1))::double precision ^ '2'::double precision) + (((a.x2 - b.x2))::double precision ^ '2'::double precision)) + (((a.x3 - b.x3))::double precision ^ '2'::double precision)) + (((a.x4 - b.x4))::double precision ^ '2'::double precision))))
   GPU Projection: a.id::integer, b.id::integer, sqrt(((((((a.x1 - b.x1))::double precision ^ '2'::double precision) + (((a.x2 - b.x2))::double precision ^ '2'::double precision)) + (((a.x3 - b.x3))::double precision ^ '2'::double precision)) + (((a.x4 - b.x4))::double precision ^ '2'::double precision)))::double precision
   Depth 1: GpuNestLoop, JoinQual: (a.id > b.id)
            Nrows (in/out: 3333333.20%), KDS-Heap (size: 8.11MB, nbatches: 1)
   Extra: bulk-exec-support
   Kernel Source: /opt/pgsql/base/pgsql_tmp/pgsql_tmp_strom_913.0.gpu
   ->  Seq Scan on public.test b  (cost=0.00..1637.00 rows=100000 width=20)
         Output: b.id, b.x1, b.x2, b.x3, b.x4
(9 rows)

少しEXPLAIN出力結果が変わっているが、GpuJoinの GPU Projection: が従前の Pseudo Scan: に相当する。
つまり、この実行計画は、GPU側で計算処理を行った上でそれをCPU側に返却。CPU側では計算結果を参照するだけ(= 計算は行わない)なので、Projection処理を省略することができるようになる。

実際のところ、もう少し適切な問題サイズの推定と分割無しには、上記のクエリのように GpuNestLoop が入力行数に対して3.3万倍の出力行を生成するようなワークロードが適切に働くとは思えないが、この辺は、後日GpuJoinの問題領域分割にDynamic Parallelismを適用する辺りでなんとかするとして、ひとまずインフラとしてはこれで良いように思える。