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を適用する辺りでなんとかするとして、ひとまずインフラとしてはこれで良いように思える。

PostgreSQLのデータ構造はなぜ並列プロセッサ向きではないか。

今年もPostgreSQL Advent Calendar 2015に参加しています。

前からちょくちょく『PG-StromってXeon Phiだとどーなんですか?』的な質問を受ける事があんですが、データ構造から見て難しいので『勘弁!』という理由を紹介してみたいと思います。

PostgreSQLのレコードは、内部的には HeapTupleHeader 構造体を先頭とする可変長データとして表現されています。

struct HeapTupleHeaderData
{
    union
    {
        HeapTupleFields t_heap;   /* MVCC関連情報 */
        DatumTupleFields t_datum; /* xmin, xmaxとか... */
    }           t_choice;
    /* current TID of this or newer tuple */
    ItemPointerData t_ctid;
    /* レコードに含まれる行数とか */
    uint16      t_infomask2;
    /* 雑多なフラグ類 */
    uint16      t_infomask;
    /*
     * ヘッダ長。ユーザデータの格納位置は
     * ((char *)htup + htup->t_hoff) からスタート
     */
    uint8       t_hoff;

    /* ここまで23 bytes */
    /* NULLビットマップ(if HEAP_HASNULL)*/
    bits8       t_bits[FLEXIBLE_ARRAY_MEMBER];

    /* MORE DATA FOLLOWS AT END OF STRUCT */
    /* この後ろにユーザ定義列の内容が詰まっている */
};

で、このレコードのxx番目の列にアクセスする、という場合は先頭から順番にたぐっていくわけです。

1番目の列はint型だからポインタを4byte進め、2番目の列はfloat型だから8byte境界までポインタを進めた上で更に8byte進め、3番目の列はNULLだからポインタは進めず、...、といった事を目的の列まで順番に進めます。

ただ、これには例外があり、

  • レコードがNULL値を含んでいない
  • 目的の列よりも前にある列に可変長データが含まれていない

場合には、当該列を格納している位置が一意に定まるので、条件を満たす場合には目的の値を1ステップで参照する事ができますが。

詳しくはNikolay Shaplov氏のTuple internals: exposing, exploring and explainingというPGconf.EU 2015での発表がよく纏まっているので、こちらを参照して頂ければと思います。


話を並列プロセッサに戻します。

現在のPG-StromはCUDA、つまりNVIDIAGPUを使うように設計されているのですが、この人は(複数のコアがプログラムポインタを共有するとはいえ)スカラプロセッサなので、各々のコアが互いに独立なメモリ領域を参照する事ができます。
つまり、CUDAコア0をレコード0にマップし、CUDAコア1をレコード1にマップし、、、、という事をすれば、各々のコアが独立にレコードの先頭から列を手繰っていけばよいだけなので、別にこれらレコードが隣接領域に配置されていなくとも処理自体は実行可能な訳です。
f:id:kaigai:20151215214346p:plain
もちろん、NVIDIAGPUメモリのバス幅は256bitとか320bitですので、複数のCUDAコアが同時に隣接領域のDRAMをアクセスすると、1回のメモリトランザクションで複数コアが使用するデータをロードできるので、もっと最適化できる・・・というのはありますが。

一方、Xeon PhiのようにSIMD命令で512bit幅の演算器を単精度×8とか倍精度×4で使う事でピーク性能を出す事を前提とするプロセッサだと、色々とデータの配置にも制約が出てきます。
f:id:kaigai:20151215214353p:plain
少なくとも、現在のPostgreSQLのデータ構造であるHeapTupleHeaderを先頭とする可変長データでレコードを表現する限り、単精度×8の演算を同時に実行できるからといって、512bit幅の演算器で8レコード分の計算を一気にこなすわけにはいきません。

これに対処する方法としては2つ考えられるのですが、

  1. CPUでデータを抽出して512bitアラインの長大配列として再配置
  2. SIMD演算器向きのデータ構造でレコードを保持する。

①はコプロセッサに処理をオフロードするためにわざわざCPUの処理が増えているのでナンセンス。
(実際、私も以前にやってみた事があったが、、、)
②は列指向データがおそらく対応する事になるとは思うのですが、PostgreSQL用のネイティブ列指向ストレージはレビュー&標準機能化に向けてもう少し時間が必要そうです。パッチ自体の規模が大きいのでちょっと二の足を踏んでしまうところではあるのですが…。

この辺、SIMD命令だけでなく、上記のようにスカラ型GPUにとってもメリットの大きい話なので、喩えて言えば『列指向ストレージが入ったら本気だす』といったところでしょうか。

Sort by Table Partition?

v9.6向け開発ネタとして思い付いたアイデア
でも、個人的には他に優先すべき機能*1もあるので、たぶん自分ではできない。誰かヨロシク的な。

タイムスタンプをキーとして複数の子テーブルにパーティション化されたテーブルがあるとする。
これは結構一般的な伝票データの作り方なのでそれほど変な仮定でもない。

各子テーブルに設定されたCHECK()制約から、特定のキーによる並べ替えを行う場合に各々の子テーブルに大小関係が定義できる場合。
例えば、以下のようなテーブル構成で、キー "YMD" によるソートを行うケースを考えると、tbl_2013テーブルに格納されている全てのレコードは、他のテーブルから読み出したレコードよりも最近の日付を持っていると言える。中を読むまでもなく。

f:id:kaigai:20150610000543p:plain

そうすると、キー "YMD" による並べ替えを行うケースであっても、ソートを行う問題領域を小さくする事ができるので、その分、処理時間を短くできる。

要は、こういった条件が満たされる場合には

GroupAggregate
 -> Sort (key: YMD)
   -> Append
     -> SeqScan on tbl_2013
     -> SeqScan on tbl_2012
     -> SeqScan on tbl_2011
     -> SeqScan on tbl_2010

これが

GroupAggregate
 -> Append
   -> Sort (key: YMD)
     -> SeqScan on tbl_2013
   -> Sort (key: YMD)
     -> SeqScan on tbl_2012
   -> Sort (key: YMD)
     -> SeqScan on tbl_2011
   -> Sort (key: YMD)
     -> SeqScan on tbl_2010

これが等価になると思う。

QuickSortの計算量は平均で O(NLogN) なので、仮に個々の子テーブルが同じ数のレコードを持っていた場合、たった4分割のケースであっても、ソートの計算量は後者の方が半分くらいで済む事になる。

もちろんプランナがもっと賢くならないとダメなんだが、Appendのパスを作る時点で『ここには後でキー"YMD"によるソートが要求される可能性がある』という事が判っていれば、pathkeys付きでAppendパスを追加してやれば良いのではなかろうか。

*1:Aggregate Before Joinとかね

NVRTCライブラリを使う

CUDA7.0RCの新機能の一つに、Runtime Compilationというのがある。

従来、cuModuleLoad()などでGPUバイナリをロードして使う際には、nvccを実行してC/C++のソースからPTXなりのバイナリを生成する必要があった。CUDA 7.0RCのRuntime Compilationは、これをOpenCLライクな関数呼び出しでソース→バイナリへの変換を行うライブラリである。
予め静的なコードを書いてコンパイルしておく、というレベルの話であれば一向に構わないのだが、PG-Stromの場合はGPUのコードは実行時に動的生成されるので、GPUコードのビルド時間は割と性能影響の大きいファクター。あまりチンタラやっている訳にはいかない。

以下の図はCUDA Toolkitのドキュメントにあるnvccの処理フローを示した図だが、nvccに何を食わせるかによっても処理フローが変わってくる。つまり、ビルド時間に影響がある。
f:id:kaigai:20150307232258p:plain

まず考えねばならないのは、入力となるソースコードは何であるか。
.cu形式と.gpu形式が選択肢ではあるが、どちらも一長一短である。

  • .cu形式の場合、先ずホストコード/デバイスコードの分離を行い、さらにC/C++で二通りのビルドプロセスが走る。要はもっさりしている。その一方で、CUDA関数が利用できるよう適切なファイルを勝手にインクルードしてくれたりと、デプロイを考える上では非常に楽。
  • .gpu形式の場合、予めデバイスコードとして分離されたものに対してビルドを行うので、ビルド時間は非常に短い。ただ、.cu形式であればnvccが勝手に解決してくれたようなデプロイにまつわるあれやこれやを自分で指定してやる必要がある。正直、これをNVIDIA謹製ツール以外が行うのは違うと思う。

という悩みがあった訳だが、CUDA 7.0RCでサポートされたNVRTCはこの辺の問題をきれいさっぱり解決してくれた。

まず最初に、cstring形式のソースコードを与えて nvrtcProgram オブジェクトというものを作成する。

nvrtcResult nvrtcCreateProgram(nvrtcProgram *prog,
                               const char *src,
                               const char *name,
                               int numHeaders,
                               const char **headers,
                               const char **includeNames)

次に、このnvrtcProgramオブジェクトをビルドする。
nvccのオプションを全てサポートしている訳ではないが、いくつかのコンパイルオプションを指定する事もできる。

nvrtcResult nvrtcCompileProgram(nvrtcProgram prog,
                                int numOptions,
                                const char **options)

で、ビルドが終わった後のnvrtcProgramオブジェクトからは、PTXイメージやビルドログを取り出す事ができる。

nvrtcResult nvrtcGetPTX(nvrtcProgram prog, char *ptx)

nvrtcResult nvrtcGetProgramLog(nvrtcProgram prog, char *log)

当初、このAPIを見た時に、nvrtcCompileProgram()関数の呼び出しが同期呼び出しになってしまうので、ソースのビルド中にテーブルスキャンを先に進められないのはイカガナモノカ…と思ったのだが、よく考えたら、こんなものはDynamic Background Workerを使えば何とでもできる。

f:id:kaigai:20150307232301p:plain
例えば、エグゼキュータの開始時に nvrtcCompileProgram を呼び出して(あるいは、従来はnvccコマンドを実行して)、GPUコードのビルドを行うDynamic Background Workerプロセスを立ち上げておく。
この時点ではまだ、GPUで命令を処理するためのバイナリは生成されていないので、PostgreSQL本体側はせっせとテーブルのスキャンに励み、データをDMAバッファに積んでおく。
で、しばらくするとBgWorker側でGPUコードのビルドが完了しバイナリが生成されるので、BgWorkerはSetLatch()を使ってバックエンド側にビルドが完了した事を通知する。
この時点で動作モードが切り替わり、今までメモリに読み込んできたデータブロックのDMA転送とGPUカーネルの起動が始まると共に、並行してスキャンを行うという形になる。

ここで、BgWorker側が nvcc でC/C++両方のホストコードを作ったり・・・などとしていると、下手すれば先にスキャンの方が終わってバックエンドを待たせてしまうというカッコ悪い事になるが、今回の NVRTC ライブラリを使うと、この辺の無用なオーバヘッドがない分、概ね1sec以下でビルドが終わってくれ、個人的には大変満足であった。

GpuNestedLoop

現時点でPG-Stromが対応しているワークロードは以下の4つ。

  • 全件探索 (GpuScan)
  • 表結合 (GpuHashJoin)
  • 集約演算 (GpuPreAgg)
  • ソート (GpuSort)

これに、GPU内の計算処理で使うデータ型や関数が対応しているかどうかで、GPUオフロードできるかどうかが決まる。
だいたいパッと思い付く程度にSQLクエリ処理で重い(CPUバウンドな)ヤツはカバレッジができてきたけれども、一つ大物が残っている。NestedLoop。

どう実装したものかと思っていたが、よいアイデアが浮かんだので備忘録代わりに。
f:id:kaigai:20150301211033p:plain

NestedLoopの場合、結合条件が単純な X=Y の形式ではないので、HashJoinやMergeJoinを用いて効率的に処理する事ができない。要はDBの中で総当たりを行う事になるので非常に重い。

今までに実装した上記の4つのロジックでは、PG-Stromは一次元的にGPUカーネルを起動していた。つまり、N個のGPUスレッドを起動する時にX軸にのみスレッドを割り当てていたのだが、X軸/Y軸をうまく使えばNestedLoopに落ちざるを得ないような表結合もうまく表現できるのではないかと考える。

イデアはこうだ。Inner-RelationとOuter-Relationからの入力をそれぞれ一次元の配列と捉える。
統計情報からの推定によりInner側は比較的行数が少ない方でY軸と置く。一方、Outer側は行数が多い方でX軸と置く。
で、一回のGPUカーネル実行で(Nx×Ny)個のGPUスレッドを起動すれば、各スレッドがそれぞれ対応する行のペアに対してNestedLoopの結合条件を評価し、マッチするペアのみを結果として取り出す事ができる。

CPUでのNestedLoopの実装は二重ループになっているので、如何せん時間がかかる。なので、普通はクエリ実行計画を見て真っ先に回避可能性を探る部分ではあるが、数千コアの同時並列実行能力の力でこういった制限も苦にならないとなれば、大きなアドバンテージになるだろう。

しかもこのGpuNestedLoopのロジックには、メモリアクセスで大きなアドバンテージを得られる可能性がある。
GPUスレッドはブロックという単位でグルーピングされ、同じブロックに属するGPUスレッド間は共有メモリを介したデータ共有が可能である。で、共有メモリはL1キャッシュと同じなので、DRAMへのアクセスに比べると非常に高速にアクセスが可能。
一方、GpuNestedLoopの処理ロジックの特性上、X軸上のインデックスが等しいスレッド、Y軸上のインデックスが等しいスレッドが複数個存在する。例えば、ブロックサイズ 32x32 (=1024) でGPUカーネルを起動した場合、X軸上のインデックスが 7 というGPUスレッドは他に32個存在しているハズである。これらのスレッドは Xindex=7 である行から同じデータを読み出すハズなので、何もDRAMアクセスを32回繰り返さなくても、1回だけで済む。わお!

NestedLoopに落ちざるを得ないJoinの条件は何も特別なモノでなく、例えば t0.X = t1.A OR t0.X = t1.B みたいな条件でも、NestedLoop以外のJoinロジックは使用不能となる。
一つ考慮しなければいけないのは、Inner-Relation側が十分に絞られておりGPUDRAMに載ってくれるサイズである必要がある、という事だが、そもそもNestedLoopで片側のサイズが数万行にもなるようなケースでは破綻していると言えるので、まぁ、実際上は問題なかろう。

並列Aggregateに向けて

PostgreSQL Advent Calendar 2014に参加しています。


数日前、SimonがPgSQL-Hackersに面白いパッチを投げてきた。

曰く、

KaiGai, David Rowley and myself have all made mention of various ways we could optimize aggregates.
Following WIP patch adds an extra function called a "combining function", that is intended to allow the user to specify a semantically correct way of breaking down an aggregate into multiple steps.

Gents, is this what you were thinking? If not...

もしかすると、PostgreSQL v9.5に入るかもしれないこの機能の背景を少し解説してみたい。

PostgreSQLにおいて集約関数(Aggregate Function)は、二種類の関数呼び出しによって実装されている。

  1. Transiton Function
    • 一レコード分のデータを読み込み、それに基づいて内部状態を更新する。
  2. Final Function
    • その時点の内部状態から、結果となるスカラー値を生成する

f:id:kaigai:20141219214544p:plain

具体的に、ビルトインのAVG(float)のケースを見てみる事にする。

集約関数 AVG(float) は以下のようにfloat8[3]型の内部状態を持ち、transition functionにはfloat8_accumが、final functionにはfloat8_avgが指定されている。

  • float8 transvalues[3]
    • transvalues[0] ... N (入力データの個数)
    • transvalues[1] ... 入力データの和
    • transvalues[2] ... 入力データの自乗の和(標準偏差・分散の算出で使用)

つまり、float8_accumが呼び出される度にtransvalues[0]をインクリメント、transvalues[1]に入力値を加算して内部状態を更新し、最後にfloat8_avgが呼び出された時に transvalues[1] / transvalues[0] を計算して返せば、めでたしめでたし float 型平均値を得る事ができる。

話を元に戻そう。上のパッチで、Simonは新たにcombined functionと呼ばれる関数を集約関数の定義に加えようとしている。
これは何をするものか?

KaiGai Koheiがcombined functionの役割を説明して曰く、

Its concept is good to me. I think, the new combined function should be responsible to take a state data type as argument and update state object of the aggregate function. In other words, combined function performs like transition function but can update state object according to the summary of multiple rows. Right?

つまり、集約関数の内部状態を更新するのに、一行一行を読むのではなく、別の集約関数の内部状態を引数として受け取り、それに基づいて自らの集約関数の内部状態を更新するという役割を果たす事が期待されている。

そうすると、何が起こるか?

上の平均値を求めるロジックを考えてもらいたい。AVG(float)関数の内部状態とは、いわば『N=136、合計値=12345』といった情報だ。しかし、必ずしもこれを一行毎に内部状態を更新せねばならない理由は、ない。
それぞれ別個に計算された『N=136、合計値=12345』という内部状態と、『N=123、合計値=23456』という内部状態を足し合わせ、『N=259、合計値=35801』という内部状態を作っても一向に問題はないはずだ。このように、複数行の結果をサマリした"部分集約"とでも呼ぶものを後で合成する事を可能にするのが combined function である。

では、これで何が嬉しいのか?

以下のようなクエリを考えてみよう

SELECT AVG(s.sales_price), p.prod_id
  FROM production p JOIN sales s ON p.prod_id = s.prod_id
  GROUP BY p.prod_id;

2つのテーブルをJOINし、その結果に対して集約関数AVG()を適用する。
通常、この処理はJOIN処理を行った後、JOIN結果に対して一行ずつ集約関数(のtransition function)を適用する事で行われる。

が、このクエリを以下のように書き換えたとしても同一の結果を得られる。
ここで、主クエリのAVG()はcombined functionにより中間結果を足し合わせる挙動を取る、サブクエリのAVG()はfinal functionを呼び出さず中間結果を返すものとする。

SELECT AVG(s.sales_price), p.prod_id
  FROM
    production p
  JOIN
    (SELECT AVG(sales_price), prod_id FROM sales GROUP BY prod_id) s
  ON p.prod_id = s.prod_id
  GROUP BY p.prod_id;

例えば、productionテーブルに百万行、salesテーブルに十億行のレコードが含まれていた場合、INNER JOIN処理は十億行を結合し、次いでGROUP BY prod_idはグループ毎に平均で1000行を集約する事になる。
一方、先にサブクエリで部分集約を作成し、これを後でJOINする場合、結合しなければならない行数は百万行にすぎない。

この方式には他の利点もある。テーブルをJOINした後に集約関数を適用するのに比べ、フラットなテーブルに集約関数を適用するというのは、領域分割による並列処理を実装しやすいというメリットがある。
そうすると、テーブルを並列スキャン+部分集約し、それを遥かに少ない行数のレコードをJOINするというシナリオが現実味を帯びてくる。

今回のPgSQL-Hackersでの議論は、こういった機能拡張の下地になる事を見込んだインフラ機能というワケである。

AWSでPG-Strom

PG-Stromを動かそうという時に、GPU自体は安価で手に入れやすい部品なのだけども、普段使いのLinuxサーバにそうそう都合よくGPUが挿さっている人はマジョリティではないかもしれない。

という事で、簡単にPG-Stromをお試しするために、AWSでのAMIイメージを作ってみました。

AMI-Idは ami-51231b50 で、GPUインスタンス(g2.x2large)を使えば簡単にトライアルできます。

以下に手順をご紹介しますが、デプロイ完了まで10分以下。こりゃ便利だわ・・・。

① Launch Instance を選択する
f:id:kaigai:20141112213815p:plain

② キーワード「strom」で検索するか、上記のami-51231b50でAMIイメージを検索する。
AMIイメージはちょくちょく更新されるので、キーワード検索を使った方が間違いがないかも。
f:id:kaigai:20141112214605p:plain

③ 続いてインスタンスタイプを選択。もちろんGPUインスタンス(g2.x2large)の一択です。
f:id:kaigai:20141112214700p:plain

④ 確認画面。本当はEBSストレージなどアタッチするのだろうけども、とりあえず一発起動するだけなので、そのまま「Launch」を選択。
f:id:kaigai:20141112214822p:plain

仮想マシンに接続するためのSSH鍵を選択します。
f:id:kaigai:20141112214851p:plain

⑥ これでデプロイ完了。あとは1~2分ほど起動を待つ。
f:id:kaigai:20141112214934p:plain

⑦ 初期化中。。。。
f:id:kaigai:20141112215039p:plain

⑧ さっきの秘密鍵を使ってSSHログイン
f:id:kaigai:20141112215325p:plain

⑨ PG-Strom有効なPostgreSQLインスタンスが起動している
f:id:kaigai:20141112215351p:plain

AWSのg2.x2largeタイプなので、GPUNVIDIAのGRID K520。
こんな感じでデバイスのプロパティを参照することができる。

postgres=# SELECT * FROM pgstrom_opencl_device_info();
 dnum | pnum |            property             |         value
------+------+---------------------------------+----------------------------
    0 |    0 | platform index                  | 1
    0 |    1 | platform profile                | FULL_PROFILE
    0 |    2 | platform version                | OpenCL 1.1 CUDA 6.5.18
    0 |    3 | platform name                   | NVIDIA CUDA
    0 |    4 | platform vendor                 | NVIDIA Corporation
    : |    : |    :                            |   :

テスト用のSQL~/pg_strom/testdb.sqlに置いてあるので、これを使って2000万行のテーブルx1個、4万行のテーブルx4個を作ると、昨日のエントリで紹介したテストテーブルを作成できる。

[ec2-user@ip-10-126-51-20 ~]$ psql -U postgres -f testdb.sql postgres

昨日のエントリで使ったGPUMaxwell世代のGTX980で、Kepler世代のGRID K520とは少し特性は異なるものの、まぁ、早くなってるから良いだろう。

PG-Stromありの場合

postgres=# EXPLAIN (ANALYZE, COSTS OFF) SELECT * FROM t0 NATURAL JOIN t1 NATURAL JOIN t2;
                                    QUERY PLAN
----------------------------------------------------------------------------------
 Custom (GpuHashJoin) (actual time=97.992..5512.238 rows=20000000 loops=1)
   hash clause 1: (t0.bid = t2.bid)
   hash clause 2: (t0.aid = t1.aid)
   Bulkload: On
   ->  Custom (GpuScan) on t0 (actual time=9.260..1220.530 rows=20000000 loops=1)
   ->  Custom (MultiHash) (actual time=43.361..43.362 rows=80000 loops=1)
         hash keys: bid
         Buckets: 46000  Batches: 1  Memory Usage: 99.99%
         ->  Seq Scan on t2 (actual time=0.009..9.551 rows=40000 loops=1)
         ->  Custom (MultiHash) (actual time=21.681..21.681 rows=40000 loops=1)
               hash keys: aid
               Buckets: 46000  Batches: 1  Memory Usage: 49.99%
               ->  Seq Scan on t1 (actual time=0.011..9.632 rows=40000 loops=1)
 Execution time: 9144.220 ms
(14 rows)

PG-Stromなしの場合

postgres=# SET pg_strom.enabled = off;
SET
postgres=# EXPLAIN (ANALYZE, COSTS OFF) SELECT * FROM t0 NATURAL JOIN t1 NATURAL JOIN t2;
                                   QUERY PLAN
---------------------------------------------------------------------------------
 Hash Join (actual time=46.932..29132.386 rows=20000000 loops=1)
   Hash Cond: (t0.aid = t1.aid)
   ->  Hash Join (actual time=23.504..17693.551 rows=20000000 loops=1)
         Hash Cond: (t0.bid = t2.bid)
         ->  Seq Scan on t0 (actual time=0.004..4851.203 rows=20000000 loops=1)
         ->  Hash (actual time=23.273..23.273 rows=40000 loops=1)
               Buckets: 65536  Batches: 1  Memory Usage: 2813kB
               ->  Seq Scan on t2 (actual time=0.006..10.589 rows=40000 loops=1)
   ->  Hash (actual time=23.256..23.256 rows=40000 loops=1)
         Buckets: 65536  Batches: 1  Memory Usage: 2813kB
         ->  Seq Scan on t1 (actual time=0.007..10.555 rows=40000 loops=1)
 Execution time: 32513.584 ms
(12 rows)