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)

PG-Stromなう

最近、方々で『GPUイイよ!GPU!』と言って回っている訳ですが、今現在、PG-Stromの開発がどんなもんじゃいというのをまとめておこうと思います。

振り返ってみると、2012年1月、最初にPG-Stromのプロトタイプを作ってみた時は、まさにPG-Strom管理下の外部テーブル(Foreign Scan)に対して非常に複雑な条件句を与えた場合、という非常に限られた条件下でGPUオフロードが効果を発揮するものでした。しかも、プログラミングに使ったのはCUDAなので、NVIDIA専用だったし。
f:id:kaigai:20141111222738p:plain

その後、紆余曲折を経て

  • FDW(Foreign Data Wrapper)を使うのはやめ、新たに Custom-Plan APIを設計した。
    • 外部テーブルでない、普通のテーブルに対してもGPUアクセラレーションを適用可能に
    • 全件探索以外のワークロードへも対応が可能に
  • CUDAを使うのはやめ、OpenCLライブラリを使用するようにした。
    • NVIDIA以外のGPUへの対応や、(効率はベストでないものの)CPU並列も可能に
  • 列指向データ構造を捨て、PostgreSQLに合った行指向データを扱えるように
  • 全件スキャンだけでなく、表結合(Join)と集約演算(Aggregate)に対応した

アーキテクチャ的にはこんな感じ。
f:id:kaigai:20141111222746p:plain
Custom-Plan APIを介して、全件スキャン(GpuScan)、表結合(GpuHashJoin)、集約演算(GpuPreAgg)を実装するCustom-Nodeがプランナ・エグゼキュータから呼び出される。その裏でOpenCLバイスを管理するバックグランドワーカが立ち上がり、これらのCustom-Nodeから自動生成されたGPUコードと処理すべきデータの組がメッセージキューを介して送信されるので、OpenCL Serverはこれを適宜ディスパッチしてGPUで処理する事ができる。

ま、以前の実装だとかなり適用可能なワークロードが限られていて、こんな感じの手厳しいコメントを頂くこともあったが、特にJoinをGPUで高速化できるってのは結構インパクトがあると思う。

以下の図はCPUによるHash-Joinのロジックと、PG-StromのGpuHashJoinを比較したもの。
CPUであれば、①Hash表を作り ②Outer側から一行取り出してHash表を探索 ③Projectionを行って左右の表の内容をマージした行を作り、次の処理へ回す。この②と③のプロセスをひたすらループする事になるワケだ。
一方、GpuHashJoinの場合、①Hash表を作る 部分は同じだけども、②と③の部分をGPUのコアに並列実行させる。つまり、数百~数千コア分の並列度が期待できる訳である。
f:id:kaigai:20141111222816p:plain

以下のグラフは、2億件 × 10万件 × 10万件 × .... と、結合すべき表の数を増やしていった時の処理時間を、標準PostgreSQL(9.5devel)と、PG-Strom有効化の場合で比較したもの。
一番得意なワークロードで比較しているので、多少、恣意的な部分はあるにせよ、3個のテーブルを結合する時の処理時間で11倍、9個のテーブルを結合する時の処理時間で28倍もの違いが出ている。
ま、この辺は実際に世の中で使われているワークロードと、GPUにオフロードできる部分をどういう風に近づけるかで、コストパフォーマンスが決まってくるんだろうけども。
f:id:kaigai:20141111222822p:plain

こんな感じでプロファイルを取ってみると、どの処理にどれくらい時間を要しているかが判る。
(この例では 2000万件 × 4万件 × 4万件 を結合してみた)

postgres=# SET pg_strom.perfmon = on;
SET
postgres=# EXPLAIN (ANALYZE, COSTS OFF)
           SELECT * FROM t0 NATURAL JOIN t1 NATURAL JOIN t2;
                                   QUERY PLAN
---------------------------------------------------------------------------------
 Custom (GpuHashJoin) (actual time=105.020..4153.333 rows=20000000 loops=1)
   hash clause 1: (t0.aid = t1.aid)
   hash clause 2: (t0.bid = t2.bid)
   Bulkload: On
   number of requests: 145
   total time for inner load: 29.73ms
   total time for outer load: 825.77ms
   total time to materialize: 1359.44ms
   average time in send-mq: 62us
   average time in recv-mq: 1372us
   max time to build kernel: 13us
   DMA send: 5759.42MB/sec, len: 4459.39MB, time: 774.28ms, count: 722
   DMA recv: 5661.55MB/sec, len: 2182.02MB, time: 385.41ms, count: 290
   proj kernel exec: total: 197.28ms, avg: 1360us, count: 145
   main kernel exec: total: 250.03ms, avg: 1724us, count: 145
   ->  Custom (GpuScan) on t0 (actual time=6.226..825.598 rows=20000000 loops=1)
         number of requests: 145
         total time to load: 787.28ms
   ->  Custom (MultiHash) (actual time=29.717..29.718 rows=80000 loops=1)
         hash keys: aid
         Buckets: 46000  Batches: 1  Memory Usage: 99.99%
         ->  Seq Scan on t1 (actual time=0.007..5.204 rows=40000 loops=1)
         ->  Custom (MultiHash) (actual time=16.106..16.106 rows=40000 loops=1)
               hash keys: bid
               Buckets: 46000  Batches: 1  Memory Usage: 49.99%
               ->  Seq Scan on t2 (actual time=0.018..6.000 rows=40000 loops=1)
 Execution time: 5017.388 ms
(27 rows)

GPUでの処理時間は、Joinそれ自身(main kernel exec)と結果の生成(proj kernel exec)で計450msくらい。
つまり計算能力のポテンシャルはこれ位はあるので、あとは、いかに効率的にCPU/GPUでデータを受渡しするかという所がこれからのテーマになってくるのかな。

次のエントリで、PG-Stromのデプロイについて紹介します。

さらばドイツ

2011年2月にドイツへ赴任し、以降2年10ヵ月間、SAP社とのアライアンス業務に携わっていたわけですが、11月末で任地での業務を終了し日本へと戻る事になりました。

振り返れば、初めてフランクフルトの空港に到着した後、一歩外へ出たら氷点下の世界(注:ドイツの2月です)で『こりゃトンデモない場所に来たもんだ』と思ったものですが、まぁ、住めば都で何とかなるもんですな。

オフィスはドイツ南西のWalldorf市(SAP本社がある)。住むには少し規模が小さな街なので、隣のHeidelberg市の旧市街のマンションを借り、1月遅れで嫁さんも日本からやってきました。
f:id:kaigai:20120331145434j:plain

お仕事的には、2011年前後というのはちょうどSAP社がSUSE LinuxベースのインメモリDB製品 "SAP HANA" をぶち上げ始めた頃で、主に技術面でSAPやSUSEとの折衝や、HANA認定取得の実作業というのが駐在中のメインのお仕事でした。

NEC SAP HANA向けアプライアンスサーバを販売開始
http://www.nec.co.jp/press/ja/1203/0901.html

その他にも、NECのHA製品CLUSTERPROで、SAP NW~HA製品間で連携するためのインターフェース認証を取得したりとか。

NEC、「SAP HA Interface Certification」認定を取得した「CLUSTERPRO X 3.1 for SAP NetWeaver」を発売
http://jpn.nec.com/press/201210/20121002_01.html

この手の "○○認証" という仕事が多かったのですが、SAPの認定プログラムはちょくちょくおかしくて、例えば SAP Linux 認証プログラムは80コア/160HTのNECサーバの認定を通すために、SAPの認定プログラムの方を手直しして認定を通したりとか。

一方、OSS活動は工数の20%~30%を充当して良いという事を条件に赴任したので、まだ中途であるSE-PostgreSQLの開発も継続する事ができた。これは会社に感謝。ただ、OSS活動に専従という訳ではないので、MLでの反応が遅れたりと、コミュニティに対して少し申し訳ない。
ドイツに居る間に新たに開発してみたのが PG-Strom というモジュールで、クリスマス休暇の間にCUDAを使ってプロトタイプの実装を行ってみた。これがもう一つ、自分にとって軸となるソフトウェアになるかもしれない。
http://kaigai.hatenablog.com/entry/20120106/1325852100

f:id:kaigai:20131201153401j:plain
これをプラハでのPGconf.EUで発表した際に、非常に大きな反響があった事は"コレは行けるんじゃないか?"という確信に繋がった。これはNECに限った事なのか分からないけども、自分の組織の外に対して何かを発表し、フィードバックを得るという機会が極端に少ない、あるいはナイーブになっているような気がする。
そうすると、自分の提案したい事が第三者の視点から見て価値ある事なのか、そうでないのか確信できないので、上司にちょっとダメ出しされただけで萎縮してしまうという事に繋がらないか。
先日、晴れて四捨五入したらオッサン年代に突入したので偉そうな事を書くと、若い人は公式非公式を問わず、自分のアイデアを社外で発表する機会を持った方がよい。そうすると、上司のダメ出しが的を得ているのか、的外れなのか、自分の判断基準を持てるようになるから。

この辺SAPが偉いなと思うのが、SCN(SAP Community Network)という仕組みを用意し、Wikiやブログを通して開発者やパートナー企業とエンドユーザが直接意見交換できる仕組みが機能している事。例えば製品の変な使い方やプロトタイプ的に作ってみた機能(当然公式にはサポート外)をオープンにして、それに対するフィードバックを受けられるようにしている。
もちろん、これはSAPのように非常に強力なパートナー網、顧客網というエコシステムを持っているからこそできる事。同じ事ができる企業は非常に限られている。
ただ、エコシステムの主催者でなくともコントリビュータとして社外の意見フィードバックを活用するというのは可能。まさにOSSコミュニティが日常的に行っている事で、だからこそイノベーションがそこにある。

仕事環境で言うとドイツ人が偉いのは、きっかり時間通りに仕事を切り上げる事。
2011~2013シーズンはBaden-Hillsカーリングクラブでプレーしていたのだけども、毎週火曜日19:30~のリーグ戦に皆きっちり集まるのがすごい。日本のカルチャーでは難しい所もあるけど、ぜひ見習いたい。目前の業務にばかり時間を費やして、別の領域にアンテナを張る時間や余暇の時間がなければ新しい発想なんて生まれないし、それが中長期的には組織を徐々に蝕んでいく事になるんじゃないかと思った。
f:id:kaigai:20111011183217j:plain

その他、生活面では嫁さんが考えて日本食を作ってくれたので、病気をする事もなくきっちり職務を遂行する事ができた。感謝。結婚してからの3年間、風邪をひく事すら無くなった。
ヨーロッパに住んでいる事で、なかなか日本からは行きにくい場所を旅行できた事は役得だけど、書き出すとキリがないのでないのでこれはまた別の機会に。

f:id:kaigai:20120331154344j:plain

日本に帰った後は、SAPから離れて再びOSSの仕事に戻ります。
自分で手を動かしてコードを書くと同時に、それをお金に換えていく方法をデザインする事が求められているので、タフな仕事だけども、一つのロールモデルとして後進に示せれば良いかなと。

それでは明日、ドイツを発ちます。