人生に影響を与えたプレゼンテーション

10年近くも技術者をやっていると、他の人のプレゼンテーションを見る機会はちょくちょくある。(別にOSS/Linuxに限った話じゃなくて)
全く印象に残らないものから、発表後にプレゼンターを捕まえて議論が始まるもの、中には自分の考え方や活動に影響を与えたプレゼンテーションというのもあるハズ。

パッと二つ思い浮かんだのでご紹介。
・・・まぁ、他の人はどんなプレゼンテーションを見て心動かされたのか知りたい、という下心もあるんですが。

How to Contribute to the Linux Kernel, and Why it Makes Economic Sense (James Bottomley, Novell; LinuxCon/Japan 2009)

なぜオープンソースへの貢献が重要で、ビジネス上も意義のある事である事なのかを説いたLinuxCon/Japan 2009でのキーノートの一節。
本来、企業が何か"価値"を顧客に届けるには、その価値の源泉であるUnique Innovationの部分と、その前提であるDelivery Support for the innovationの部分に投資をしなければならないが、オープンソースを活用する事で企業はUnique Innovationの部分に注力でき、コミュニティと歩調を合わせて進む(勝手forkをしない)事でDelivery Support for the innovationの部分をShared Cost(そう、freeではなくshared costなのです)に保つ事ができるという趣旨の発表。

確かに自分のケースを振り返ってみても、SE-PostgreSQLを開発するために、その価値の源泉であるSELinuxとの連携モジュールはともかくとして、SQL構文解析やクエリ最適化、クエリ実行といったRDBMSそれ自身を作るなんてのはあり得ない話。確かに自分も"巨人の肩に乗って"開発しているんだと納得したものだ。

企業がオープンソースの開発に貢献した方が、より低いコストでイノベーションを達成できる事、そして"Free"でなく"Shared cost"である事。この考え方は正に自分の人生に影響を与えたと思う。

Parallel Image Searching Using PostgreSQL and PgOpenCL (Tim Child, 3DMashUp; PGcon2011)

もう一つは、昨年のPGconで聴講したTim ChildのPG/OpenCLの発表。
内容はRDBMSに格納した画像データ(レコード長がかなり大きい)をGPUで高速に処理するというもの。
何が心に刺さったかというと、(一字一句正確ではないかもしれないが)最後のまとめで彼が語った『Database records are ideal data structure for parallel processing(データベースのレコードは理想的な並列処理のデータ構造だ)』という一言。

これを聞いて、なるほどGPUを使ってCPUをオフロードすれば、検索処理を高速に処理できるではないか・・・?という発想でCUDAやOpenCLを調査し、その半年後にできたのがPG-Stromというモジュールだったりする。

さて、他の人に聞いてみたら、どんな『自分の人生に影響を与えたプレゼンテーション』が出てくるのだろう?

PG-Stromにプロファイラをつけてみた

1月6日(金)に書いた『しゅとろ〜む、しゅとろ〜む』の記事は割と反響が大きかったようだ。
コメント欄に次のような質問を頂いたので、試してみることにする。

通りすがりさん wrote:
すばらしい成果ですね.
カラム指向的にデータを持っていること自体が性能向上に寄与しているということはないですか?
(通常 + CPU) vs (カラム指向+GPU)で比較をされていますが,
(通常 + CPU) vs (カラム指向+CPU) vs (カラム指向+GPU) の評価にも興味があります.

plan.c 内の is_device_executable_qual() 関数が常に false を返すようにすれば、条件句の処理をCPUだけで行うようになる。これは (カラム指向+CPU) と同等である。

1,000万件のレコードを持つ、通常のテーブル t1 と、PG-Strom管理下のテーブル t2 に対してそれぞれ以下のクエリを実行してみた。

■ 1回目(バッファにデータが乗っていない状態)

(通常 + CPU)
Timing is on.
postgres=# SELECT COUNT(*) FROM t1 WHERE sqrt((x-25.6)^2 + (y-12.8)^2) < 15;
 count
-------
  6718
(1 row)

Time: 8041.237 ms

(カラム指向 + CPU)
postgres=# SELECT COUNT(*) FROM t2 WHERE sqrt((x-25.6)^2 + (y-12.8)^2) < 15;
 count
-------
  6718
(1 row)

Time: 8660.486 ms

(カラム指向 + GPU)
postgres=# SELECT COUNT(*) FROM t2 WHERE sqrt((x-25.6)^2 + (y-12.8)^2) < 15;
 count
-------
  6718
(1 row)

Time: 4667.643 ms

■ 2回目(バッファにデータが乗っている状態)

(通常 + CPU)
postgres=# SELECT COUNT(*) FROM t1 WHERE sqrt((x-25.6)^2 + (y-12.8)^2) < 15;
 count
-------
  6718
(1 row)

Time: 7016.732 ms

(カラム指向 + CPU)
postgres=# SELECT COUNT(*) FROM t2 WHERE sqrt((x-25.6)^2 + (y-12.8)^2) < 15;
 count
-------
  6718
(1 row)

Time: 6733.771 ms

(カラム指向 + GPU)
postgres=# SELECT COUNT(*) FROM t2 WHERE sqrt((x-25.6)^2 + (y-12.8)^2) < 15;
 count
-------
  6718
(1 row)

Time: 173.351 ms

(通常+CPU)と(カラム指向+CPU)の比較で、ディスクからの読み出しが発生する場合にカラム指向の方が8%程度遅いという結果になっている。
複雑な条件句を設定したために、I/OよりもCPUバウンドな処理になっている事、xとyにはランダムな値を入れているために、全く圧縮が効いていないのが一因かもしれない。
(カラム指向 + GPU)で2回目の方が早くなっているのは、主にGPUコードのJITコンパイルの処理時間の違いによるものだろう。JITコンパイルにここまで時間がかかることは稀だが、確実にI/Oを発生させるために Linux の Page Cache をクリアしてから測定を行ったため、nvccコマンドもOSのキャッシュから弾き出されたという事だろう。

ただ、psql の \timing ではトータルの実行時間を表示するだけで、何が要因で時間を食っているのかは分からない。PG-Stromは性能改善を目的とするモジュールなので、どの辺を改善したら良いのか探るには先ず、どの辺にボトルネックがあるのかを探る必要がある。

という訳で、PG-StromのGUCパラメータ pg_strom.exec_profile を追加してみた。
これに "on" をセットすると、各々コンポーネントで消費した時間を表示してくれる。

postgres=# SET pg_strom.exec_profile = ON;
SET

(カラム指向 + GPU; 1回目)

postgres=# SELECT COUNT(*) FROM t2 WHERE sqrt((x-25.6)^2 + (y-12.8)^2) < 15;
INFO:  PG-Strom Exec Profile on "t2"
INFO:  Total PG-Strom consumed time: 4367.067 ms
INFO:  Time to JIT Compile GPU code: 1741.505 ms
INFO:  Time to initialize devices:   345.353 ms
INFO:  Time to Load column-stores:   2119.669 ms
INFO:  Time to Scan column-stores:   3.566 ms
INFO:  Time to Fetch virtual tuples: 110.920 ms
INFO:  Time of GPU Synchronization:  31.244 ms
INFO:  Time of Async memcpy:         31.320 ms
INFO:  Time of Async kernel exec:    27.906 ms
INFO:  Num of registers/thread [0]:  25
INFO:  Constant memory usage [0]:    40 byte
INFO:  Max device memory usage[0]:   536 KB
 count
-------
  6718
(1 row)

Time: 4514.738 ms

\timing で計測した応答時間 4514.738ms のうち、PG-Strom モジュール内の処理時間は 4367.067 msで、そのうち、大部分を占めるのが、GPUコードのJITコンパイル(1741.505ms)と、カラムストアからのロード(2119.669ms)になる。これと比べると、GPUでの処理時間・メモリ転送は桁が違う。

(カラム指向 + GPU; 2回目)

postgres=# SELECT COUNT(*) FROM t2 WHERE sqrt((x-25.6)^2 + (y-12.8)^2) < 15;
INFO:  PG-Strom Exec Profile on "t2"
INFO:  Total PG-Strom consumed time: 183.302 ms
INFO:  Time to JIT Compile GPU code: 0.043 ms
INFO:  Time to initialize devices:   1.134 ms
INFO:  Time to Load column-stores:   54.883 ms
INFO:  Time to Scan column-stores:   3.425 ms
INFO:  Time to Fetch virtual tuples: 96.384 ms
INFO:  Time of GPU Synchronization:  27.462 ms
INFO:  Time of Async memcpy:         30.737 ms
INFO:  Time of Async kernel exec:    27.906 ms
INFO:  Num of registers/thread [0]:  25
INFO:  Constant memory usage [0]:    40 byte
INFO:  Max device memory usage[0]:   536 KB
 count
-------
  6718
(1 row)

Time: 186.867 ms

1回目で時間を食っていた、GPUコードのJITコンパイル処理時間が消え、カラムストアからのロード時間も大幅に減っている。また、地味にデバイスの初期化にも345.353 ms要していたが、これがほぼ無くなっている。
この結果、トータルの処理時間が4514.738 ms⇒186.867msに減少。
カラムストアのロード/スキャンと、タプルをフェッチする処理(これはFDWの仕様なので減らすのが難しい)、それにGPUの処理の同期で合わせて 182.154 ms が消費されている。

1/6(金)の時点から少しアルゴリズムを変更しているが、メモリ使用量はほとんど問題になっていない。
これは、I/O周りで時間がかかっているために、2個、3個とチャンクを非同期に処理しようとしても、次のチャンクを読み込んでGPUに渡す頃には、前のチャンクの処理が既に終わっているからという事だろう。
この辺、もっと足回りの良いマシンなら変わってくるのだろうか。

なお、Time to scan... というのは、条件句を評価した結果に基づいてカラムストアをスキャンする処理で、条件句には使われていないものの、Target-listに含まれるカラムが存在する場合に発生する。今回のクエリは COUNT(*) を返すだけなので、追加のスキャンは発生していない。


おまけ。(カラム指向 + CPU)の実行結果だとこうなる。

postgres=# SELECT COUNT(*) FROM t2 WHERE sqrt((x-25.6)^2 + (y-12.8)^2) < 15;
INFO:  PG-Strom Exec Profile on "t2"
INFO:  Total PG-Strom consumed time: 2314.374 ms
INFO:  Time to JIT Compile GPU code: 0.000 ms
INFO:  Time to initialize devices:   0.000 ms
INFO:  Time to Load column-stores:   6.881 ms
INFO:  Time to Scan column-stores:   1435.570 ms
INFO:  Time to Fetch virtual tuples: 871.891 ms
INFO:  Time of GPU Synchronization:  0.000 ms
INFO:  Time of Async memcpy:         0.000 ms
INFO:  Time of Async kernel exec:    0.000 ms
 count
-------
  6718
(1 row)

Time: 8063.461 ms

トータル 8063ms のうち、PS-Strom内の処理は 2314 ms。つまり、必死こいてPG-Stromから本体側にメモリコピーの後、CPUで条件句を処理という流れが見える。PG-Strom内での結果の絞込みができないので、Fetch virtual tuplesの時間が大幅に増加しているのが分かる。
それと、Scan column-store の時間もやや気がかり。足回りとして、この辺は改善の余地があるやも。

PG-Strom

I've checked up an idea whether it is feasible to implement, or not, since I saw a presentation by Tim Child in Ottawa last year.
Is it possible to accelerate sequential-scan of PostgreSQL?
We often see sequential-scan instead of index-scan in case of queries with complex calculation. I thought GPU works fine in these cases.
I tried to implement a module that works as FDW (foreign data wrapper) of PostgreSQL, since I could have a time to develop during Christmas vacation.

The name of module is PG-Strom that is pronounced as shutt-row-me; being pronounced in German style.
Its name originates "Streaming Multiprocessor" that is a unit of process in GPU.

Of course, it assumes existing interface of FDW, so it is unavailable to update, and some more restrictions like sort or aggregate functions. However, it achieves good performance as a prototype.

Note that the following description is based on author's understanding (quite newbie for CUDA), so please point out if something incorrect.

Benchmark

Even though it is an arbitrary testcase, I tries to execute a query that scans a table with 20-million records in my development environment. NVidia's GTS450eco is installed.

-- A regular table
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

-- with 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

It is a surprising result. PG-Strom returns the result with 10 times faster!
In addition, we may be able to expect more improvement because GPU is quite cheap one (about 100Euro).

Let's try again. I reduced the number of records (5-million records, with shared_buffer=960MB) to store whole of the table on the buffer; to eliminate affects from disk-I/O.

-- A regular table
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

-- with PG-Strom
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

Wow!

Idea

PostgreSQL iterates (1) fetch a tuple from storage (or buffer), and (2) evaluation of qualifier of WHERE clause according to contents of the tuple during sequential-scan. Thus, it unavailable to handle (2) during execution of (1), and also unavailable to handle (1) during execution of (2). An idea is CPU multi-threading, however, it is hard to implement because PostgreSQL does not have thread-safe design including memory or I/O management.

PG-Strom entrusts GPU device the (2) portion (evaluation of WHERE clause), and make CPU focus on I/O stuff.
The calculation stuff shall be handled on GPU device side asynchronously, so it shall be finished during CPU handles more I/O stuff.

However, GPU is not a magic bullet for anything.

We need to transfer data to be calculated by GPU into device memory mounted on GPU. It requires to transfer via PCI-E that has narrow bandwidth compared to the one between CPU and Memory. (Max 2.5GB/s in x16 lane)
Thus, amount of data to be copied should be smaller as we can as possible.

In most cases, it is rare case that WHERE clause reference all the columns within the table, because the purpose of query is to fetch a record that satisfies the condition of XXXXX.
PG-Strom handles execution of WHERE clause on GPU device. At that time, all copied to GPU device are contents of referenced columns. I expect 10%-20% of table size needs to be copied to GPU device via PCI-E, because numeric data is smaller than text data.

Data structure and Asynchronous process

The internal data structure of PG-Strom is organized according to the above idea.
For example, when we create a foreign table with four-columns: a, b, c and d, PG-Strom creates tables corresponding to each columns within pg_strom schema. These tables have rowid (int64) to identify a particular row and an array-type to store multiple original data.

Even though it is a column-oriented data structure recently well used, it does not go out of transaction management of PostgreSQL, PG-Strom does not need to touch them.

This type of data structure allows PG-Strom to load data into GPU devices via PCI-E bus effectively.
The contents read from the databases are temporarily stored on fixed-length buffer called "chunk", then it shall be moved to GPU devices and calculated, and the results shall be written back at last. These steps are executed asynchronously, thus, CPU can scan the database concurrently to set up next chunk. This design enables to utilize both of CPU and GPU.

Just-in-time compile and native-code execution

CPU and GPU have its own advantage and disadvantage for each. GPU has much higher computing capability using large number of calculation units in parallel, however, one of its disadvantage is conditional branch.

NVidia's GPU synchronously run 32 of execution units (that is called as SM:Streaming-Multiprocessor) like as a SIMD operations. In the case when device code contains conditional-branch part, a particular thread has 'true' on the condition, and other thread has 'false' on the condition, then, all the threads execute both of true-block and false-block and result of the block to be skipped shall be ignored. Thus, we cannot ignore the cost to execute branch statement within GPU device, especially, if-block is big.

PostgreSQL has internal representation of WHERE clause as tree-structure, and we scan the tree-structure using switch statement on execute them. It shall be worst effectiveness.

Thus, PG-Strom adopts Just-in-time compile to generate native binary code of GPU to avoid execution control on GPU device.

When the supplied query tries to reference a foreign-table managed by PG-Strom, the query planner requires PG-Strom to generate execution plan. At that time, PG-Strom dynamically generate a source code towards GPU device, then kicks nvcc (compiler of NVidia's device) to build a native code of GPU device.
Of course, it shall be cached on shared memory to avoid execute compiler so frequently.

Next, when query-executor calls PG-Strom's executor, as I mentioned above, this native code shall be transferred to the device side with data read from pg_strom schema, and executed asynchronously.
The qualifiers of WHERE clause is already extracted on the planner stage, no need to handle a big switch statement.

We can confirm the automatically generated code of GPU device.

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)

Publication

Right now, it is in public at GitHub. GPLv3 is applied.
https://github.com/kaigai/pg_strom

Even though it is a prototype, thus, its specification depends on my feeling, and we cannot expect documentation for a while, if you'd like to use, please call me (@kkaigai) on twitter.

A short demonstration

This is a short demonstration. The 't1' table is a regular table with 5-million records, and the 't2' table is a foreign table managed by PG-Strom also with 5-million records.
In the case of sequential-scan with complex qualifier, scan on 't2' was finished x10 times faster than the case of 't1'.

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

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

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

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

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

ベンチマーク

かなり恣意的なテストケースではあるが、2,000万件のレコードからなるテーブルを全件スキャンするクエリを手元の環境で実施してみた。なお、搭載しているGPUNvidia 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は非常に多くの並列演算ユニットを協調して動作させる事により高い計算能力を発揮するが、不得意な分野もある。その一つが条件分岐である。

NVidiaGPUでは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(NVidiaGPU向けコンパイラ)を実行して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) などで呼びかけてもらえれば。

Leaky Views と Security Barrier : PostgreSQL Advent Calendar #4

このエントリはPostgreSQL Advent Calendarに参加しています。12/4(日)担当也。ヨーロッパ中部時間ではまだ12/3(土)ですが。



RDBMSで行レベルのアクセス制御を実現する方法として、利用者に対して直接のアクセス権を付与せずに、特定のビューを通してだけアクセスを許可するのはしばしば使われるテクニックです。
ですが、場合によっては不可視な行の中身を参照できてしまうというのは、あまり広く認知されている訳ではないようです。
ので、問題のポイントと、現在開発中の PostgreSQL v9.2 に提案しているアイデアをご紹介します。

ユーザ定義関数のCOST値による問題

ここでは、以下の表を例に考えてみます。

customerテーブル

列名 制約
cid int primary key
cname text not null
cmail text
cpasswd text

customerテーブルには全顧客の情報が格納されているため、利用者は自分自身の情報しか見る事ができないよう設定しましょう。
(便宜上 PostgreSQL ユーザ名が cname に対応するものとします)

postgres=# CREATE VIEW my_account AS SELECT * FROM customer
                  WHERE cname = getpgusername();
CREATE VIEW
postgres=# GRANT SELECT ON my_account TO public;
GRANT

本来、このテーブルには 3ユーザ分の情報が格納されているのですが、

postgres=# SELECT * FROM customer;
 cid | cname |       cmail       | cpasswd
-----+-------+-------------------+----------
 101 | alice | alice@example.com | abcdef
 102 | bob   | bob@example.com   | xyz123
 103 | eve   | eve@example.com   | deadbeaf
(3 rows)

確かに、自分自身の情報しか参照できないように見えます。

postgres=# SET SESSION AUTHORIZATION alice;
SET
postgres=> SELECT * FROM customer;
ERROR:  permission denied for relation customer
postgres=> SELECT * FROM my_account;
 cid | cname |       cmail       | cpasswd
-----+-------+-------------------+---------
 101 | alice | alice@example.com | abcdef
(1 row)

しかし、利用者がSQL関数を定義できる場合、面白い事が起こります。
publicスキーマはデフォルトでCREATE権限を全体に与えている事に注意!)

postgres=> CREATE FUNCTION f_leak(text) RETURNS bool LANGUAGE plpgsql
           COST 0.00000001
           AS 'BEGIN RAISE NOTICE ''f_leak => %'', $1; RETURN true; END';
CREATE FUNCTION

postgres=> SELECT * FROM my_account WHERE f_leak(cmail);
NOTICE:  f_leak => alice@example.com
NOTICE:  f_leak => bob@example.com
NOTICE:  f_leak => eve@example.com
 cid | cname |       cmail       | cpasswd
-----+-------+-------------------+---------
 101 | alice | alice@example.com | abcdef
(1 row)

おっと、何か見えてはならないモノが見えたようです。
どういう事なのでしょうか、ちょっと EXPLAIN で調べてみましょう。

postgres=> EXPLAIN SELECT * FROM my_account WHERE f_leak(cmail);
                           QUERY PLAN
-----------------------------------------------------------------
 Seq Scan on customer  (cost=0.00..20.85 rows=1 width=100)
   Filter: (f_leak(cmail) AND (cname = (getpgusername())::text))
(2 rows)

この実行計画はVIEWの本体である customer テーブルをスキャンしていますが、利用者が付与した f_leak() とVIEWの条件を順にチェックしています。
問題は、副作用を持つ f_leak() の実行コストが非常に小さな値に設定されているため、オプティマイザは cname = getpgusername() より先にf_leak()を実行して不必要な条件の判断を省略した方が得策であると判断して、関数の実行順序を並べ替えている事です。その結果、不可視であるべき行の内容が引数としてf_leak()に渡され、それが利用者に漏えいしている訳です。

JOINと条件句の分配に伴う問題

同様に、VIEWによる行レベルアクセス制御を破るシナリオはもう一つ知られています。
先ほどの customer テーブルに加えて、もう一つテーブルを追加して考察を進めてみましょう。

creditテーブル

列名 制約
cid int references customer(cid)
number text
expired date

この credit テーブルは顧客のクレジットカード番号を保持しています。先ほどの my_account ビューと同様に、自分自身のレコードだけを参照できるようなVIEWを定義してみましょう。

postgres=# SELECT * FROM customer;
 cid | cname |       cmail       | cpasswd
-----+-------+-------------------+----------
 101 | alice | alice@example.com | abcdef
 102 | bob   | bob@example.com   | xyz123
 103 | eve   | eve@example.com   | deadbeaf
(3 rows)

postgres=# SELECT * FROM credit;
 cid |       number        |  expired
-----+---------------------+------------
 101 | 1111-2222-3333-4444 | 2014-02-28
 102 | 5555-6666-7777-8888 | 2013-10-30
 102 | 1234-5678-1234-5678 | 2015-06-30
 103 | 0987-6543-2109-8765 | 2014-08-31
(4 rows)

postgres=# CREATE VIEW my_credit AS SELECT cname, cmail, credit.*
           FROM customer NATURAL JOIN credit WHERE cname = getpgusername();
CREATE VIEW
postgres=# GRANT SELECT ON my_credit TO public;
GRANT

おや、やっぱり何かおかしいようです。

postgres=# SET SESSION AUTHORIZATION alice;
SET
postgres=> SELECT * FROM my_credit;
 cname |       cmail       | cid |       number        |  expired
-------+-------------------+-----+---------------------+------------
 alice | alice@example.com | 101 | 1111-2222-3333-4444 | 2014-02-28
(1 row)

postgres=> SELECT * FROM my_credit WHERE f_leak(number);
NOTICE:  f_leak => 1111-2222-3333-4444
NOTICE:  f_leak => 5555-6666-7777-8888
NOTICE:  f_leak => 1234-5678-1234-5678
NOTICE:  f_leak => 0987-6543-2109-8765
 cname |       cmail       | cid |       number        |  expired
-------+-------------------+-----+---------------------+------------
 alice | alice@example.com | 101 | 1111-2222-3333-4444 | 2014-02-28
(1 row)

もう一度 EXPLAIN で実行計画を眺めてみましょう。

postgres=> EXPLAIN SELECT * FROM my_credit WHERE f_leak(number);
                              QUERY PLAN
----------------------------------------------------------------------
 Hash Join  (cost=20.89..43.96 rows=2 width=104)
   Hash Cond: (credit.cid = customer.cid)
   ->  Seq Scan on credit  (cost=0.00..21.60 rows=387 width=40)
         Filter: f_leak(number)
   ->  Hash  (cost=20.85..20.85 rows=3 width=68)
         ->  Seq Scan on customer  (cost=0.00..20.85 rows=3 width=68)
               Filter: (cname = (getpgusername())::text)
(7 rows)

困ったことに、『creditテーブルをf_leak()条件でスキャンした結果』と『customerテーブルをcname = getpgusername()条件でスキャンした結果』がJOINされています。
オプティマイザはJOINすべき行を最小化するよう条件句を分配するのですが、f_leak()関数は credit テーブルの number 列のみ、cname = getpgusername() 条件は customer テーブルの cname 列のみに依存しています。そのため、JOINの完了を待つ事なく個々のテーブルをスキャンする時点で条件句を実行した方が、JOINすべき行数を減らす事ができます。
その結果、副作用を持つf_leak()がcreditテーブルのスキャン計画に push-down され、最初の例と同様に、不可視であるべき行の内容がf_leak()に渡され、それが利用者に漏えいしてしまっています。

この2つの問題は共に、オプティマイザがVIEWの境界を越えて関数の実行順序を入れ替えている事が原因です。これは性能観点からは優れた実装ですが、セキュリティを目的としたVIEW定義という観点では問題です。
一方で、VIEW内部で使われている関数を全て評価してから、その外部から与えられた関数を評価するという実装は、安全ですが、性能上無視できない性能劣化をもたらします。例えば、1万行 x 1万行のテーブルをJOINする場合、外部から与えられた関数をテーブルスキャンの時点で評価する事で片方の行数を1万行から100行に絞り込めるとしたら、9900万行分のJOIN処理を省略する事ができます。

次に、PostgreSQL v9.2に向けて提案されている Leaky View 問題への対策を紹介しましょう。

VIEW の security_barrier 属性と最適化の抑制

ここからは、私の提案している「Fix Leaky View Problemパッチ」の解説です。
前節で考察したように、VIEWを行レベルアクセス制御の目的で利用する場合には、パフォーマンスとセキュリティのトレードオフが存在します。安全側に倒せば許容できない程の性能劣化を招く可能性があり、一方、性能最適であれば情報漏えいの危険があります。

Fix Leaky Views Problem パッチは、CREATE VIEW構文でWITH(...)句を用いてオプション値を指定することを許容します。構文は以下の通りです。

CREATE VIEW view_name [WITH (options[,...])] AS select_statement;
options:
  security_barrier[= true|false]

security_barrier オプションは、VIEWが行レベルアクセス制御を目的として定義されていることを示す属性です。これを指定することで、一部のクエリ最適化を抑制する事が可能になります。
この設計に至るまでには長い議論があったのですが、結局、パフォーマンスとセキュリティのどちらが重要であるのかを判断できるのはVIEWを定義する人のみである、というシンプルな結論にたどり着いたのでした。
VIEWにsecurity_barrier属性が付与されている時、VIEWの内側で使用されている全ての関数・条件句は、VIEWの外側から与えられた関数・条件句よりも先に実行される事が保証されます。

では、実際に試してみましょう。以下で定義する my_account_secure と my_credit_secure は、先ほどの2つの例で使用したVIEWにsecurity_barrier属性を付加したものです。

postgres=# CREATE VIEW my_credit_secure WITH (security_barrier) AS
           SELECT cname, cmail, credit.* FROM customer NATURAL JOIN credit
           WHERE cname = getpgusername();
CREATE VIEW
postgres=# GRANT SELECT ON my_account_secure TO public;
GRANT
postgres=# CREATE VIEW my_account_secure WITH (security_barrier) AS
           SELECT * FROM customer WHERE cname = getpgusername();
CREATE VIEW
postgres=# GRANT SELECT ON my_credit_secure TO public;
GRANT

動作結果は以下のようになりました。"f_leak => ..." と表示されている内容は、クエリによって本来参照可能なデータの範囲内に収まっている事が分かります。

postgres=# SET SESSION AUTHORIZATION alice;
SET
postgres=> SELECT * FROM my_account_secure WHERE f_leak(cmail);
NOTICE:  f_leak => alice@example.com
 cid | cname |       cmail       | cpasswd
-----+-------+-------------------+---------
 101 | alice | alice@example.com | abcdef
(1 row)

postgres=> SELECT * FROM my_credit_secure WHERE f_leak(number);
NOTICE:  f_leak => 1111-2222-3333-4444
 cname |       cmail       | cid |       number        |  expired
-------+-------------------+-----+---------------------+------------
 alice | alice@example.com | 101 | 1111-2222-3333-4444 | 2014-02-28
(1 row)

では、VIEWにsecurity_barrier属性を付加することで、クエリ実行計画にどのように変化しているのでしょうか?先ほどと同じように、EXPLAIN構文で調べてみましょう。

postgres=> EXPLAIN SELECT * FROM my_account_secure WHERE f_leak(cmail);
                               QUERY PLAN
-------------------------------------------------------------------------
 Subquery Scan on my_account_secure  (cost=0.00..20.88 rows=1 width=100)
   Filter: f_leak(my_account_secure.cmail)
   ->  Seq Scan on customer  (cost=0.00..20.85 rows=3 width=100)
         Filter: (cname = (getpgusername())::text)
(4 rows)

f_leak()関数の評価は cname = getpgusername() 条件で customer テーブルをスキャンした後に行われる事が分かります。オプティマイザは security_viwe 属性を持ったVIEWの内側に条件句を push-down しなくなりました。

もう一つの例も同様です。

postgres=> EXPLAIN SELECT * FROM my_credit_secure WHERE f_leak(cmail);
                                 QUERY PLAN
----------------------------------------------------------------------------
 Subquery Scan on my_credit_secure  (cost=20.89..46.96 rows=2 width=104)
   Filter: f_leak(my_credit_secure.cmail)
   ->  Hash Join  (cost=20.89..46.90 rows=6 width=104)
         Hash Cond: (credit.cid = customer.cid)
         ->  Seq Scan on credit  (cost=0.00..21.60 rows=1160 width=40)
         ->  Hash  (cost=20.85..20.85 rows=3 width=68)
               ->  Seq Scan on customer  (cost=0.00..20.85 rows=3 width=68)
                     Filter: (cname = (getpgusername())::text)
(8 rows)

パッチ自体の動作原理は極めて単純です。

PostgreSQLは、一旦、VIEWに対するクエリを内部的にサブクエリに書き換えます。その後、オプティマイザがクエリ実行計画を作成する際に、"シンプル"なサブクエリ(OFFSET/LIMIT句を含まない…など)であれば、性能最適の観点からサブクエリをJOINを用いてフラット化(Pull-Up)します。
その後で、条件句はオプティマイザによって性能上最適な位置に振り分けられるため、VIEWの内側/外側といった区別はもはや意味を持たなくなります。

VIEWのsecurity_barrier属性は、この際の条件に作用します。RangeTblEntry構造体のsecurity_barrierは、関連するサブクエリがVIEWに由来し、かつ、VIEWのsecurity_barrier属性がtrueである場合にセットされます。
以下の処理では、security_barrier属性が false だとpull_up_simple_subquery()は呼ばれないため、サブクエリのフラット化は抑制されます。

--- a/src/backend/optimizer/prep/prepjointree.c
+++ b/src/backend/optimizer/prep/prepjointree.c
@@ -543,6 +543,7 @@ pull_up_subqueries(PlannerInfo *root, Node *jtnode,
         */
        if (rte->rtekind == RTE_SUBQUERY &&
            is_simple_subquery(rte->subquery) &&
+           !rte->security_barrier &&
            (containing_appendrel == NULL ||
             is_safe_append_member(rte->subquery)))
            return pull_up_simple_subquery(root, jtnode, rte,

さらにもう一ヶ所。条件句に与える引数が特定のサブクエリにだけ依存している場合、オプティマイザはこの条件句の実行をサブクエリ処理の中に移動(Push-Down)しようとしますが、同様にサブクエリが security_view 属性つきのVIEWに由来する時は、これをスキップします。

@@ -763,6 +769,7 @@ set_subquery_pathlist(PlannerInfo *root, RelOptInfo *rel,
      Node       *clause = (Node *) rinfo->clause;

      if (!rinfo->pseudoconstant &&
+         !rte->security_barrier &&
          qual_is_pushdown_safe(subquery, rti, clause, differentTypes))
      {
          /* Push it down */

この2ヶ所の処理を追加することによって、これまで見たような、VIEWを行レベルアクセス制御の目的に使用する場合の問題を回避する事ができます。

FUNCTION の leakproof 属性

Leaky View問題はVIEWのsecurity_barrier属性によって解決する事ができるのですが、これは一部のクエリ最適化を無効化するために、場合によっては、そのためのコストが看過できないほど大きい事もあります。

例えば、アプリケーションの設計上、以下のようなVIEWを定義し、VIEWの外側から条件句(主キーによる絞込みなど)を与えて使いたいというケースを考えてみましょう。

CREATE VIEW valid_credit WITH (security_barrier) AS
    SELECT * FROM credit WHERE card_is_valid(number, expired);

SELECT * FROM valid_credit WHERE cid = <customer-id>

この場合、card_is_valid関数と、VIEWの外部から与えた cid = <customer-id> を用いて credit テーブルをスキャンした結果が利用者には返されます。ですが、VIEWにはsecurity_barrier属性が設定されているため、常にcard_is_valid関数が先に実行されます。
この制限は cid 列にインデックスが設定されていても同様です。したがってインデックス・スキャンが選択されるべき状況でも全件スキャンが選択されてしまいます。ああ困った、困った。

Fix Leaky View ProblemパッチはPart-1とPart-2から構成されており、Part-1は前述の security_barrier 属性の実装を、Part-2ではその例外を設定する機能を実装しています。

Part-2によって、先ほどのオプティマイザへの変更は一部修正されます。

--- a/src/backend/optimizer/path/allpaths.c
+++ b/src/backend/optimizer/path/allpaths.c
@@ -769,7 +769,8 @@ set_subquery_pathlist(PlannerInfo *root, RelOptInfo *rel,
        Node       *clause = (Node *) rinfo->clause;

        if (!rinfo->pseudoconstant &&
-           !rte->security_barrier &&
+           (!rte->security_barrier ||
+            !contain_leakable_functions(clause)) &&
            qual_is_pushdown_safe(subquery, rti, clause, differentTypes))
        {
            /* Push it down */

サブクエリがsecurity_barrier属性付きのVIEWに由来するとき、このif文は条件句のPush-Downを抑止しますが、Part-2パッチは条件句(clause)が leakable-functions (つまり情報を漏えいする可能性のある関数)を含んでいなければ、サブクエリへの条件句のPush-Downを許可するように修正します。

では、関数が情報を漏えいする可能性の有無をどのように設定するか。
それには、CREATE FUNCTION構文に新たに追加されるLEAKPROOF属性を使用します。

例えば、以下のように使用します。LEAKPROOFを指定することで、この関数に情報漏えいの恐れがないという事を明示的に指定できますが、これは同時に、潜在的に不可視の行の内容を参照することを可能にするため、関数のLEAKPROOF属性をセットするには特権ユーザの権限が必要です。
SE-PostgreSQLでも db_procedure:{install}権限をチェックする予定です)

CREATE FUNCTION is_positive(int) RETURNS bool LANGUAGE plpgsql
    LEAKPROOF
    AS 'BEGIN RETURN $1 > 0; END';

一部のビルトイン関数の中でも、明らかに情報漏えいのリスクがない関数については、デフォルトでLEAKPROOF属性がセットされています。
(全部で2400個程あるため、網羅的なチェックはこれからですが…。)

例えば、32bit Integer同士の大小比較を行う int4gt 関数は、以下のように実装されています。

Datum
int4gt(PG_FUNCTION_ARGS)
{
    int32       arg1 = PG_GETARG_INT32(0);
    int32       arg2 = PG_GETARG_INT32(1);

    PG_RETURN_BOOL(arg1 > arg2);
}

この実装に情報漏えいの危険はありませんので、DB初期化時にLEAKPROOF属性はセットされています。
その他にも、現在のパッチでは各種ビルトインタイプの等価・大小比較演算子の実装として利用されている関数にLEAKPROOF属性がついています。実際に試してみましょう。

postgres=# SET SESSION AUTHORIZATION bob;
SET
postgres=> SELECT * FROM my_credit;
 cname |      cmail      | cid |       number        |  expired
-------+-----------------+-----+---------------------+------------
 bob   | bob@example.com | 102 | 5555-6666-7777-8888 | 2013-10-30
 bob   | bob@example.com | 102 | 1234-5678-1234-5678 | 2015-06-30
(2 rows)

ユーザ bob は2枚のクレジットカードを持っています。リッチメンですね。

では、2つの条件句を付加してみます。一つは先ほどのf_leak()関数、もう一つは expired < '2014-01-01' という Date 型の大小比較演算です。

postgres=> SELECT * FROM my_credit_secure WHERE f_leak(number) AND expired < '2014-01-01';
NOTICE:  f_leak => 5555-6666-7777-8888
 cname |      cmail      | cid |       number        |  expired
-------+-----------------+-----+---------------------+------------
 bob   | bob@example.com | 102 | 5555-6666-7777-8888 | 2013-10-30
(1 row)

NOTICEメッセージが一行だけ表示されているという事は、大小比較演算はf_leak()関数よりも先に実行されたようです。EXPLAINで実行計画を見てみましょう。

postgres=> EXPLAIN SELECT * FROM my_credit_secure WHERE f_leak(number) AND expired < '2014-01-01';
                                QUERY PLAN
---------------------------------------------------------------------------
 Subquery Scan on my_credit_secure  (cost=1.06..27.06 rows=1 width=104)
   Filter: f_leak(my_credit_secure.number)
   ->  Hash Join  (cost=1.06..27.04 rows=2 width=104)
         Hash Cond: (credit.cid = customer.cid)
         ->  Seq Scan on credit  (cost=0.00..24.50 rows=387 width=40)
               Filter: (expired < '2014-01-01'::date)
         ->  Hash  (cost=1.05..1.05 rows=1 width=68)
               ->  Seq Scan on customer  (cost=0.00..1.05 rows=1 width=68)
                     Filter: (cname = (getpgusername())::text)
(9 rows)

見ての通り、expired < '2014-01-01' 条件句が credit テーブルのスキャンに結びついているのと比較して、f_leak()関数はmy_credit_secureビューの内側にPush-Downされていません。これが LEAKPROOF 属性の有無による違いです。もし credit テーブルにインデックスが設定されていれば、Push-Downされた条件句により、全件スキャンの代わりにインデックス・スキャンが選択されるかもしれません。

まとめ

確かこの問題は、かれこれ2年以上議論を続けてきた息の長い問題です。

2009年9月4日のセキュアOS塾SE-PostgreSQL vs Oracle Label Security』の資料の中で言及があります。(p.34)
http://sepgsql.googlecode.com/files/090904-jsosjk04-sepgsql-vs-ols.pdf

開発コミュニティとしての方向性は、概ね上で紹介した形で収束しつつありますが、まだ v9.2 の新機能として紹介できるかどうか、は分からない状況です。が、SE-PostgreSQLの行レベルアクセス制御機能を実現するためにもマージしておきたい機能ですので、なんとかcommitできるよう頑張りたいところです。

最後に『じゃあ、既存のシステムではどうやって対策したら良いのよ?』という質問に対して一つTIPSを紹介しておきたいと思います。

Q. PostgreSQL v9.1以前のバージョンでLeaky View問題を防ぐにはどうしたらよいか?
A. クエリに OFFSET 0 を付ける

オプティマイザがサブクエリをフラット化、または、条件句をPush-Downする時、サブクエリにOFFSET/LIMIT句が含まれている場合はそれを断念する、という事を思い出してください。
OFFSET 0は結果セットの先頭から値を読むという意味ですので、本来は何の意味もありません。ですが、ここまで説明した条件句の実行順序に起因する問題を防ぐには簡便な方法です。
ただし、関数のLEAKPROOF属性に相当する機能はありませんので、その点でトレードオフは必要になります。



PostgreSQL Advent Calendar向けに記事を書くにあたり、MySQL、MS SQL ServerOracle Databaseなど他のRDBMSの挙動はどうなっているのか調べたかったのですが、時間がありませんでした。特に Oracle は勝手にWHERE句に条件をくっつけるVirtual Private Databaseという機能を持っていますので気になります。

これらは、追って調査したいと思います。きっと。いつの日か。アディオス、アミーゴ。


さて、翌 12/5(月) は笠原さんです。よろしく〜

Leaky VIEW まとめ

SELinuxとは関係のない、RDBMSでのセキュリティのお話。

利用者に対して、テーブルに対する直接のアクセス権を与えず、特定のビューを通してだけアクセスを許可するのは、行レベルのアクセス制御でよく使われるテクニックである。
つまり、ビューは不可視であるタプルをフィルタリングする役割を持つ。

しかし、これで万全かというと、そうではない。
クエリ最適化を上手く利用することで、利用者が見えないはずのタプルを参照する事は可能である。

以下の例を見て頂きたい。

postgres=# CREATE TABLE T1 (id int, name text);
CREATE TABLE
postgres=# CREATE TABLE T2 (id int, cred text);
CREATE TABLE
postgres=# INSERT INTO t1 VALUES (1, 'coke'), (2, 'soda'),
                                 (3, 'juice'), (4, 'fanta');
INSERT 0 4
postgres=# INSERT INTO t2 VALUES (1, 'public'), (2, 'hidden'),
                                 (3, 'hidden'), (4, 'public');
INSERT 0 4
postgres=# CREATE VIEW v1 AS SELECT * FROM t1 NATURAL JOIN t2
                                      WHERE t2.cred = 'public';
CREATE VIEW
postgres=# SELECT * FROM v1;
 id | name  |  cred
----+-------+--------
  1 | coke  | public
  4 | fanta | public
(2 rows)

postgres=# GRANT SELECT ON v1 TO alice;
GRANT

ビュー v1 は、テーブル t1 と t2 を JOINし、ここでは t2.cred = 'public' が行レベルのセキュリティポリシー、すなわち、フィルタリング対象の行を定めるものとする。

ユーザ alice は t1 と t2 へのアクセス権を持っていないため、ビュー v1 を通してしか、これらの情報にアクセスできないはずである。

だがしかし、以下のクエリの実行結果を見てもらいたい。

postgres=> SELECT getpgusername();
 getpgusername
---------------
 alice
(1 row)

postgres=> SELECT * FROM t1;
ERROR:  permission denied for relation t1
postgres=> SELECT * FROM t2;
ERROR:  permission denied for relation t2
postgres=> SELECT * FROM v1;
 id | name  |  cred
----+-------+--------
  1 | coke  | public
  4 | fanta | public
(2 rows)

上記の結果は想定通りだろう。
では、続いて、WHERE句にユーザ定義関数を付加する。
この関数は、常に true を返すが、引数を利用者のコンソールに出力する。

postgres=> CREATE OR REPLACE FUNCTION f_leak(text)
               RETURNS bool LANGUAGE 'plpgsql'
               AS 'BEGIN
                       raise notice ''f_lead: (%)'', $1;
                       RETURN true;
                   END';
CREATE FUNCTION
postgres=> SELECT * FROM v1 WHERE f_leak(name);
NOTICE:  f_lead: (coke)
NOTICE:  f_lead: (soda)
NOTICE:  f_lead: (juice)
NOTICE:  f_lead: (fanta)
 id | name  |  cred
                                        • -
1 | coke | public 4 | fanta | public (2 rows)

結果セットは2行だが、利用者コンソールには見えてはならないはずのタプルの内容が出力されている。

その理由はEXPLAIN分の出力を見ると明らかである。

postgres=> EXPLAIN SELECT * FROM v1 WHERE f_leak(name);
                           QUERY PLAN
                                                                                                                              • -
Hash Join (cost=25.45..356.91 rows=12 width=68) Hash Cond: (t1.id = t2.id) -> Seq Scan on t1 (cost=0.00..329.80 rows=410 width=36) Filter: f_leak(name) -> Hash (cost=25.38..25.38 rows=6 width=36) -> Seq Scan on t2 (cost=0.00..25.38 rows=6 width=36) Filter: (cred = 'public'::text) (7 rows)

f_leak()関数を探してみると、Join-Loopの内側で t1 テーブルを読み出す際のフィルタリング条件として実行されている事がわかる。

これは、f_leak()の引数が t1 由来のデータだけを参照しているため、Joinすべき行数を減らすために、本来実行されるべき位置(t1.id = t2.id を評価した後)からオプティマイザによって移動させられた事による。

とは言え、この手の最適化を行わなければビューを介したアクセスは極端に性能が悪くなるはずなので、問題は PostgreSQL に限った話ではないと思われる。

例えば、100万件のタプルを持つテーブルでID列にインデックスが張られており、処理コストの比較的高い f_policy() 関数によってフィルタリングを行うビューを介してアクセスするとする。その場合、ビューの外から ID = 1234 という条件が来た場合に常に全件スキャンが走るようなら、泣ける。

手元にOracleの環境がある友人に試してもらったところ、同様に、見えないはずのタプルの内容を出力できるそうな。

なお、PostgreSQLには、セキュリティポリシーの適用されている t2 の内容を見る方法もある。*1

以下のような関数を定義する。ポイントは COST=0.0001 の部分。

postgres=> CREATE OR REPLACE FUNCTION f_leak(text)
               RETURNS bool LANGUAGE 'plpgsql'
               COST 0.0001
               AS 'BEGIN
                       raise notice ''f_lead: (%)'', $1;
                       RETURN true;
                   END';
CREATE FUNCTION

今度は f_leak() 関数で t2 の情報を参照するようにすると、同様にフィルタリングされているはずの行の内容が出力される。

postgres=> SELECT * FROM v1 WHERE f_leak(cred);
NOTICE:  f_lead: (public)
NOTICE:  f_lead: (hidden)
NOTICE:  f_lead: (hidden)
NOTICE:  f_lead: (public)
 id | name  |  cred
                                        • -
1 | coke | public 4 | fanta | public (2 rows)

EXPLAIN文の結果

postgres=> EXPLAIN SELECT * FROM v1 WHERE f_leak(cred);
                            QUERY PLAN
                                                                                                                                  • -
Hash Join (cost=25.40..52.43 rows=12 width=68) Hash Cond: (t1.id = t2.id) -> Seq Scan on t1 (cost=0.00..22.30 rows=1230 width=36) -> Hash (cost=25.38..25.38 rows=2 width=36) -> Seq Scan on t2 (cost=0.00..25.38 rows=2 width=36) Filter: (f_leak(cred) AND (cred = 'public'::text)) (6 rows)

今度は、f_leak()がt2のScan-Loopに結合されているが、注目すべきはその順序。

f_leak()のコスト値を低く設定したために、複数のフィルタリング条件がScan-Loopに結合している場合、f_leak()が cred = 'public' よりも優先されている。

後者のシナリオは、何もJoinを伴わない場合でも実行可能である。

この問題は、既に開発者の中では既知の問題で、利用者から特定のタプルを不可視にする目的でビューを使うべきでない事が明記されている。
http://www.postgresql.jp/document/current/html/rules-privileges.html

5/18にオタワで開催される PostgreSQL Developer Meeting では、この問題を議論するつもりである。
一応、解決策の腹案は持っているが、そこまで踏み込まないにしても、先ずはこの辺のシナリオが『解決すべき課題である』というコンセンサス形成あたりを目標としたい。

*1:PostgreSQL固有というのは、他のRDBMSでユーザ定義関数のコストを指定する手段があるかどうか不明のため。

OSS開発勉強会-09

今日の勉強会で説明した内容のメモ

GitHubなんかを使って、何か別のOSSプロジェクト(gitで管理されている)の
派生バージョンを作り、かつ、オリジナルの更新に追従する方法。

オリジナルのGitリポジトリのURLを git://github.com/kaigai/my_oss.git とし、
このリポジトリに対するコミット権は無いものとする。

自分の派生バージョンのGitリポジトリのURLを ssh://git@github.com/kaigai/modified.git とする。
自分はこのリポジトリの所有者であるのでコミット権があるものとする。

% git clone ssh://git@github.com/kaigai/modified.git
% git remote add upstream git://github.com/kaigai/my_oss.git
  • 手順3:オリジナルの master をトラッキングするブランチの作成
% git checkout --track upstream/master -b develop

ここでは develop というブランチを作成し、同時に、このブランチが upstream リポジトリの master ブランチをトラッキング(追跡)するように設定している。

これにより、オリジナル側で新たなパッチが commit された場合には、git pull でオリジナルの修正をマージすることができるようになる。

  • 手順4:オリジナル側の更新をマージ
% git pull [--rebase]
  • 手順5:origin のリモートリポジトリに登録
% git push origin develop

この操作により、developブランチに加えた全ての修正がoriginのリモートブランチに加えられる事になり、他の利用者にも visible となる。

gitの場合、ローカル環境で commit しても、pushしない限りその修正はローカルのコピー(git cloneした)に留まることに留意。

この辺の git の機能を使えば、非常に楽に派生バージョンを並行に開発することができる。
SE-PostgreSQLのように、マージまでに非常に長い時間がかかって
いる場合でも、本家の追従を容易に行うことができる。