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関数を起動する事ができるようになる。

事前に問題サイズが明らかな場合にはそれほど嬉しい機能ではないのだが、問題は、Kernel関数を複数のステップに分けて実行しなければならないケースでかつ、問題サイズが前ステップの結果に依存する、すなわち、事前に問題サイズを正確に予測する事が難しいケースで非常に役に立つ。

典型例が、複数段から成るJOIN。

上記の例は、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を行うため、これも実際の入力値となる行数は事前には分からない。
一般的に、不確定要素の上に不確定要素を積み重ねていくほど誤差は大きくなり、クエリ実行計画の最適化というのもこういった誤差との戦いである。閑話休題。

これに対処するには２通りの方法が考えられる。

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() 関数に与えてロードしていたパスが少し変わってくる。

といっても、それほど大変な違いではなく、PTX形式のデータと外部のライブラリをリンクするために cuLinkXXXX() 関数を呼び出し、cubinというまた別の実行形式フォーマットを生成して、最終的に同じ cuModuleLoadData() 関数へロードする。

で、ここまで頑張った結果のパフォーマンスは以下のような結果となった。

いつものように、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ありで、縦軸は応答速度。つまり、短いほどベター。

その他の条件に関しては以下の通り。

• CPU: Intel Xeon E5-2670 v3 (12core, 2.3GHz) x2
• GPU: NVIDIA GTX980 (2048core, 1.1GHz) x1
• RAM: 384GB
• OS: Red Hat Enterprise Linux 6.6
• CUDA: CUDA Toolkit 7.5 + Nvidia driver 352.68

グラフから見て分かるように、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が非対応になる事を意味しており、現時点でインパクトが大きいのはAWSのGPUインスタンス*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がぶん回っているのが分かる。

TargetListの計算をGPUで行う。

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

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

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

-- 任意の二点間の４次元空間上の距離が 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)

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

-- 任意の二点間の４次元空間上の距離を算出して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、つまりNVIDIAのGPUを使うように設計されているのですが、この人は（複数のコアがプログラムポインタを共有するとはいえ）スカラプロセッサなので、各々のコアが互いに独立なメモリ領域を参照する事ができます。
つまり、CUDAコア0をレコード0にマップし、CUDAコア1をレコード1にマップし、、、、という事をすれば、各々のコアが独立にレコードの先頭から列を手繰っていけばよいだけなので、別にこれらレコードが隣接領域に配置されていなくとも処理自体は実行可能な訳です。

もちろん、NVIDIAのGPUメモリのバス幅は256bitとか320bitですので、複数のCUDAコアが同時に隣接領域のDRAMをアクセスすると、１回のメモリトランザクションで複数コアが使用するデータをロードできるので、もっと最適化できる・・・というのはありますが。

一方、Xeon PhiのようにSIMD命令で512bit幅の演算器を単精度×8とか倍精度×4で使う事でピーク性能を出す事を前提とするプロセッサだと、色々とデータの配置にも制約が出てきます。

少なくとも、現在のPostgreSQLのデータ構造であるHeapTupleHeaderを先頭とする可変長データでレコードを表現する限り、単精度×8の演算を同時に実行できるからといって、512bit幅の演算器で8レコード分の計算を一気にこなすわけにはいきません。

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

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

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

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

Sort by Table Partition?

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

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

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

そうすると、キー "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) なので、仮に個々の子テーブルが同じ数のレコードを持っていた場合、たった４分割のケースであっても、ソートの計算量は後者の方が半分くらいで済む事になる。

もちろんプランナがもっと賢くならないとダメなんだが、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に何を食わせるかによっても処理フローが変わってくる。つまり、ビルド時間に影響がある。

まず考えねばならないのは、入力となるソースコードは何であるか。
.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を使えば何とでもできる。

例えば、エグゼキュータの開始時に 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。

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

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側が十分に絞られておりGPUのDRAMに載ってくれるサイズである必要がある、という事だが、そもそもNestedLoopで片側のサイズが数万行にもなるようなケースでは破綻していると言えるので、まぁ、実際上は問題なかろう。