しゅとろ〜む、しゅとろ〜む

昨年、オタワでTim Child氏の発表を聞いて以来、実装できないものかと思って暖めていたアイデアがある。GPUの処理能力を使って、PostgreSQLの検索処理を高速化できないか？というものである。
特に複雑な計算を含むクエリの場合、Index-Scanに落ちないで、全件スキャンが走ることが往々にしてあるが、こういったケースで有効に作用するのではなかろうか？という着想である。
クリスマス休暇の間、割とまとまった開発時間を取る事ができたので、PostgreSQLのFDW(Foreign Data Wrapper)として動作するモジュールを作成してみた。

モジュールの名前は PG-Strom で、ドイツ風に『しゅとろ〜む』と発音する。
これは GPU の処理単位である Streaming Multiprocessor に由来する。

もちろん、現状のFDWのI/F前提なので、更新は不可能でソートや集約関数もモジュール側に出せないという諸々制約はあるが、プロトタイプとしてはまずまずの性能である。

※ なお、下記のGPU関連の記述は著者（CUDAプログラミング歴１ヶ月）の理解によるものです。間違っていたらご指摘ください。むしろ教えてくださいｗ

ベンチマーク

かなり恣意的なテストケースではあるが、2,000万件のレコードからなるテーブルを全件スキャンするクエリを手元の環境で実施してみた。なお、搭載しているGPUはNvidia GTX450ecoである。

-- 従来のテーブル
mytest=# SELECT count(*) FROM pgbench_accounts
    WHERE (xval - 23.45) * (xval - 23.45) + (yval - 54.32) * (yval - 54.32) < 100;
 count
--------
 629291
(1 row)

Time: 29030.738 ms

-- PG-Stromを利用
mytest=# SELECT count(*) FROM pgstrom_accounts
    WHERE (xval - 23.45) * (xval - 23.45) + (yval - 54.32) * (yval - 54.32) < 100;
 count
--------
 629291
(1 row)

Time: 2337.475 ms

驚いた事に、1/10以下の応答時間でクエリを実行してしまったではないか。
しかも利用しているGPUは100Euro程度のショボイものだけに、伸びしろもあるだろう。

もう一回、今度はディスクIOの影響を除くため、テーブル全体がバッファに乗るサイズ（shared_buffer=960MBで、件数を500万件に削減）で試してみた。

-- 従来のテーブル
mytest=# SELECT count(*) FROM t1
   WHERE (xval - 23.45) * (xval - 23.45) + (yval - 54.32) * (yval - 54.32) < 100;
 count
--------
 157800
(1 row)

Time: 4106.045 ms

mytest=# SELECT count(*) FROM t2
    WHERE (xval - 23.45) * (xval - 23.45) + (yval - 54.32) * (yval - 54.32) < 100;
 count
--------
 157800
(1 row)

Time: 393.346 ms

わお！

アイデア

PostgreSQLの場合、基本的に全件スキャン時の処理は (1) ディスク(or バッファ)からの読み出し (2) タプルの内容に基づいて WHERE 条件句を評価 の繰り返しとなる。
そのため、(1)の処理中は(2)を実行できず、(2)の処理中は(1)を実行できない。CPUマルチスレッド化はひとつのアイデアだが、PostgreSQLはメモリ管理やI/O周りを含めて Thread-Safe な構造にはなっていないため、これは非常に難しい。

PG-Stromでは、(2)のWHERE条件句の処理を GPU 側に任せる事で、CPUをI/Oに専念させる。
計算処理はGPU側で非同期に実施してくれるので、CPUから見た場合『ここにあるデータを評価しておいて頂戴』と頼んでおくと、しばらくI/O処理をしている間に計算結果が出来上がっている、という算段である。

ただ、GPUに処理をさせれば万事OKかというと、そうは問屋が卸さない。

GPUで計算させるには、GPU搭載のdevice memoryにデータを転送する必要があるが、これには PCI-Eを通して転送する必要があり、この箇所の帯域はCPU-Memory間に比べて非常に小さいのである。(x16のバスでもMAX片側2.5GB/s)
したがって、GPUデバイスに転送するデータの量はできるだけ少なくした方がよい。

通常、SQLのWHERE条件句がテーブルの全てのカラムを参照するという事は考えにくい。
『○○の条件を満たすレコードを取り出す』というのがクエリの目的だからだ。
PG-StromではWHERE条件句の処理をGPU側で実行するが、その際、GPUデバイス側に転送されるのは計算に必要なカラムだけである。普通は数値データの方が文字列よりも短いため、PCI-Eを介してGPUデバイスに転送の必要があるのは、テーブル全体の10%-20%程度ではなかろうか。

データ構造と非同期処理

PG-Stromの内部データ構造も、上記の方針に従って編成されている。
例えば、a、b、c、dの4つのカラムを持つFOREIGN TABLEを定義したとき、PG-Stromは各々のカラムに対応するテーブルを"pg_strom"スキーマ内に作成する。これらのテーブルは、行を一意に識別する rowid (int64) と、元々のデータを配列化したデータ型を持つ事になる。

最近流行のカラム指向DB的なデータ構造という訳だが、あくまでも PostgreSQL のトランザクション管理の枠内でデータ構造を規定しているので、その辺の厄介な処理は PG-Strom の側では一切ノータッチで済ませている。

この様なデータ構造を持つ事により、PG-StromではPCI-Eを介してGPUデバイスに送り込むデータを高速にDBから読み込めるようになっている。読み込んだデータはチャンクと呼ぶ固定長のバッファに蓄えられ、順次GPUデバイスに送出、GPUでの演算処理を行い、結果の書き戻しが行われる。
実際にはこれらの一連の処理は全て非同期に実施されるため、CPUはその間もDBからデータを読み込み、次のチャンクのセットアップが可能であるため、CPU/GPUを効率的に利用する事ができる。

Just-in-time compile と native code 実行

CPUとGPUにはそれぞれ得意不得意の分野があり、GPUは非常に多くの並列演算ユニットを協調して動作させる事により高い計算能力を発揮するが、不得意な分野もある。その一つが条件分岐である。

NVidiaのGPUでは32個の実行ユニットを含むStreaming Multiprocessorという単位で、SIMDライクな処理が行われる。GPU内の処理が条件分岐を含み、特定のスレッドでは IF 条件が真に、別のスレッドでは偽になるような場合、全てのスレッドがIF文の真ブロック/偽ブロックを処理し、IF条件に合致しないケースを破棄するという処理が行われる。そのため、特にIFブロックのサイズが大きくなるに従って、GPU内で条件句を処理する際のコストが無視できないものとなる。

PostgreSQL内部ではWHERE条件句をツリー状のデータ構造によって保持しているが、ツリーを順にスキャンして『次は '+' 演算子だから…』と switch() 文で分岐させるような処理は、最悪の効率、という事になる。

※ ただ、並列に実行する全てのスレッドでIF条件の評価結果が同じ場合にどうなるか？という点は、調べた限りではよく分からなかった。この場合にペナルティが避けられるのであれば、GPU側でコントロール処理を行うのも一つのアイデア。

そのため、PG-StromではJust-in-time compileの技術を使って動的にネイティブのGPUコードを生成して実行するという方針を採用している。

利用者のクエリがPG-Strom管理下の外部テーブルを参照する場合、Query PlannerはPG-Stromに対してクエリ実行プランを作成するよう要求する。その時、PG-Strom PlannerはWHERE条件句に従って動的にGPUデバイス用のコードを生成し、nvcc（NVidia の GPU向けコンパイラ）を実行してGPU向けのネイティブコードを生成する。もちろん、毎回コンパイラを起動していては時間の無駄なので、生成したバイナリは共有メモリ上にキャッシュされる。

次いで、Query-ExecutorがPG-Strom Executorを呼び出すと、前述の通り、pg_stromスキーマ内から読み出したデータと共に、GPU向けのネイティブコードがデバイス側に送出され、非同期に実行される。
WHERE条件句は既にPlanner段階で展開されているので、改めて巨大な switch 文を処理する必要は…ない。

ちなみに、EXPLAIN文でどのようなGPU向けのコードが生成されているかを見る事ができる。

mytest=# EXPLAIN SELECT * FROM pgstrom_accounts
         WHERE (xval - 23.45) * (xval - 23.45) +
               (yval - 54.32) * (yval - 54.32) < 100;
                        QUERY PLAN
--------------------------------------------------------------
 Foreign Scan on pgstrom_accounts  (cost=2.00..0.00 rows=1000 width=368)
    Required Cols : aid, bid, abalance, filler, xval, yval
   Used in clause : xval, yval
      1: typedef unsigned long size_t;
      2: typedef long __clock_t;
      3: typedef __clock_t clock_t;
      4: #include "crt/device_runtime.h"
      5:
      6: typedef char  bool_t;
      7:
      8: __global__ void
      9: pgstrom_qual(unsigned char rowmap[],
     10:              double c5_values[],
     11:              unsigned char c5_nulls[],
     12:              double c6_values[],
     13:              unsigned char c6_nulls[])
     14: {
     15:     int offset_base = blockIdx.x * blockDim.x + threadIdx.x;
     16:     int offset = offset_base * 8;
     17:     unsigned char result = rowmap[offset_base];
     18:     unsigned char errors = 0;
     19:     unsigned char cn5 = c5_nulls[offset_base];
     20:     unsigned char cn6 = c6_nulls[offset_base];
     21:     int bitmask;
     22:
     23:     for (bitmask=1; bitmask < 256; bitmask <<= 1)
     24:     {
     25:         double cv5 = c5_values[offset];
     26:         double cv6 = c6_values[offset];
     27:
     28:         if ((result & bitmask) &&
                    !((((cv5 - 23.45) * (cv5 - 23.45)) +
                       ((cv6 - 54.32) * (cv6 - 54.32))) < 100))
     29:             result &= ~bitmask;
     30:         offset++;
     31:     }
     32:     rowmap[offset_base] = (result & ~errors);
     33: }
(36 rows)

公開先

今のところGitHUBで公開中。ライセンスはGPLv3です。
https://github.com/kaigai/pg_strom

まだプロトタイプ段階なので、私の気分次第で仕様は変わりますし、当面はドキュメントも期待できません。それでも使ってみようという奇特な方がいらっしゃいましたら、Twitter (@kkaigai) などで呼びかけてもらえれば。