SE-PostgreSQLでリモートプロセスのセキュリティラベルを使用する

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

何を設定する必要があるのか?

何か世のため人のためになりそうなネタは・・・という事で、SE-PostgreSQLでリモートプロセスのセキュリティラベルを使用する方法をまとめてみました。

v9.1からサポートされたSE-PostgreSQLは、OSの機能であるSELinuxと連携して、SQLを介したDBオブジェクトへのアクセスに対して強制アクセス制御を行う機能です。
かいつまんで言えば、○○のテーブルは機密度の高い情報を持っている、□□は機密度が低いと言った形で、PostgreSQL標準のDBアクセス制御機能に加えて、OSのセキュリティポリシーに従ってアクセス制御を行います。そのため、同じ”情報”であるにも関わらず、保存先がファイルかデータベースかといった”格納手段”によって異なるポリシーでアクセス制御が行われる事を避ける事ができます。
SE-PostgreSQLが追加のアクセス制御を行う際には、『誰が(subject)』『何に(object)』『何をする(action)』のかをひとまとめにして、OSの一部であるSELinuxに伺いを立てるのですが、この時に『誰が』『何に』を識別するために、セキュリティコンテキストと呼ばれる特別な識別子が使用されます。
セキュリティコンテキストというのはこんな感じです。ファイルのセキュリティコンテキストを確認するには ls -Z を、プロセスの場合は ps -Z や id -Z を実行して確認できます。

[kaigai@iwashi ~]$ ls -Z
drwxr-xr-x. kaigai users unconfined_u:object_r:user_home_t:s0 patch/
drwxr-xr-x. kaigai users unconfined_u:object_r:user_home_t:s0 repo/
drwxr-xr-x. kaigai users unconfined_u:object_r:user_home_t:s0 rpmbuild/
drwxr-xr-x. kaigai users unconfined_u:object_r:user_home_t:s0 source/
drwxr-xr-x. kaigai users unconfined_u:object_r:user_home_t:s0 tmp/
[kaigai@iwashi ~]$ ps -Z
LABEL                             PID TTY          TIME CMD
unconfined_u:unconfined_r:unconfined_t:s0-s0:c0.c1023 28791 pts/4 00:00:00 bash
unconfined_u:unconfined_r:unconfined_t:s0-s0:c0.c1023 29093 pts/4 00:00:00 ps

DBオブジェクトの場合はSECURITY LABEL構文を用いて設定できるほか、システムビューのpg_seclabelsを用いて参照する事ができます。

postgres=# SELECT provider,label FROM pg_seclabels
              WHERE objtype = 'table' AND objname = 't1';
 provider |                  label
----------+------------------------------------------
 selinux  | unconfined_u:object_r:sepgsql_table_t:s0
(1 row)

そしてDB利用者の場合。やっと今回の記事の主題にたどり着けたワケですが、SE-PostgreSQLは『接続元プロセスのセキュリティコンテキスト』をDB利用者のセキュリティコンテキスト(= DB利用者の権限)として扱います。

SELinuxgetpeercon(3)というAPIを持っており、この人はソケットのファイルディスクリプタを引数として与えると、接続相手のセキュリティコンテキストを返します。SE-PostgreSQLはこのAPIを使用して『接続元プロセスのセキュリティコンテキスト』を取得しています。
で、UNIXドメインソケットによるローカル接続の場合は特別な設定は何も必要ないのですが、TCP/IPによるリモート接続の場合、OSは相手側プロセスの情報を持っていないため、ネットワーク接続の確立に先立って双方のセキュリティコンテキストを交換する Labeled Network の設定を行う必要があります。
この機能、Linux kernel 2.6.16 からサポートされた機能で、IPsecの鍵交換デーモンである racoon と連携してネットワーク接続先とセキュリティコンテキストを交換します。つまり、IPsecで通信経路が保護されている(少なくとも、改ざん検知できる)事が前提です。IPsecが前提ですよ。大切な事なので2回書きました。

Linux kernelのバージョンからも分かるようにそれほど新しい機能ではないのですが、意外と設定手順がドキュメント化されていませんので、少しまとめてみました。

テスト環境の構成

IPsecの設定を極める事が目的ではありませんので、今回は、最も単純なHost-to-HostのPSK(Pre-Shared Key)の構成でIPsec通信経路を作成します。その上で、Labeled Networkの設定を追加する事にします。


今回のテスト環境では、2台のホストをそれぞれサーバとクライアントに見立てて設定を行いました。kg1(192.168.1.42)がクライアント、kg2(192.168.1.43)がサーバであると想定し、この2台の間にIPsecの接続を設定します。ディストリビューションFedora 17 を使用し、Development Workstation構成でクリーンインストールしました。

なお、IPsec接続の設定に関してはFedora Projectのドキュメントを参考にしています。
詳しくはIPsec Host-to-Host Configurationを参照してください。

事前の設定

パッケージの追加

インストール直後の状態では ipsec-tools と system-config-network パッケージが導入されていませんので、両ホストでこれらのパッケージをインストールします。

# yum install -y ipsec-tools system-config-network

また、サーバ側には postgresql-server と postgresql-contrib パッケージが、クライアント側には postgresql パッケージが必要です。

# yum install -y postgresql-server postgresql-contrib
# yum install -y postgresql
ファイアウォールの設定

IPsecの鍵交換デーモンの通信とPostgreSQLサーバへの接続にはファイアウォールの設定変更が必要です。
system-config-firewallを実行して関連するポートを解放してください。

IPsecの鍵交換デーモン用ポートは、サーバ側/クライアント側の双方で解放が必要です。


PostgreSQLサーバ用ポートは、サーバ側で解放が必要です。

これらの設定を保存し、有効化してください。

NetworkManagerの無効化

ちょっとダサいのですが、NetworkManager経由でうまくipsec0デバイスを自動起動させる事ができなかったので、レガシーなnetworkスクリプトによってネットワークデバイスをup/downさせる事にします。

# chkconfig NetworkManager off
# chkconfig network on

Host-to-Host IPsec デバイスの作成

今回は、system-config-networkを使ってipsec0デバイスの定義ファイルや、racoonの設定ファイルを自動生成させる方法でHost-to-HostのIPsecデバイスを作成します。
まず、system-config-networkを実行して設定GUIを起動します。

起動画面です。『IPsec』のタブを選択します。


左上の『New』をクリックすると、ウィザード形式でIPsec設定を追加することができます。


デバイス名を入力します。ここでは『ipsec0』としました。


今回は『Host-to-Host encryption』を選択します。


『Automatic encryption mode selection via IKA (racoon)』を選択します。


この設定での接続相手先のIPアドレスを入力します。画面は192.168.1.42ホストのものなので、接続先は『192.168.1.43』となります。


事前共有型の認証キーを入力します。ここでは『sepgsql920』としました。


入力したパラメータを確認して『Apply』をクリックします。


これで『ipsec0』デバイスが作成されているのが分かります。


設定を保存すると、右上の『Activate』『Deactivate』を選択することができるようになりますので、『Activate』をクリックしてください。

もう一方のノードでも同様の設定を行います。

IPsecの動作確認

ひとまず、SELinux関連の設定が絡まない範囲でHost-to-Host設定のIPsecがちゃんと動作しているかを確認します。双方のノードでネットワークを再起動します。

# service network restart

tcpdumpでパケットを観察すると、IPsecのAH/ESPプロトコルが使われている事が分かります。

[root@kg2 ~]# tcpdump -n -i eth1 host 192.168.1.42
tcpdump: verbose output suppressed, use -v or -vv for full protocol decode
listening on eth1, link-type EN10MB (Ethernet), capture size 65535 bytes
16:55:34.979277 IP 192.168.1.42 > 192.168.1.43: \
    AH(spi=0x03850f37,seq=0x10): ESP(spi=0x0f3c888b,seq=0x10), length 76
16:55:34.979505 IP 192.168.1.43 > 192.168.1.42: \
    AH(spi=0x0ebad42b,seq=0x6): ESP(spi=0x0183d073,seq=0x6), length 76
16:55:34.980000 IP 192.168.1.42 > 192.168.1.43: \
    AH(spi=0x03850f37,seq=0x11): ESP(spi=0x0f3c888b,seq=0x11), length 68
    :
    :

Labeled Network の設定

ipsec0デバイスを起動する時に、どの相手との接続にIPsecを使うのか、その際のアルゴリズムは何なのか…という事がOS側に伝えられます。これを行うのがsetkeyコマンドなのですが、通常、これはifcfg-ipsec0の記述に基づいて(ここにスクリプトファイル名)が自動的に実行されます。
実際には以下のような書式をsetkeyコマンドに与えるのですが、Labeled Networkを使用するにあたって重要なのは、2行目の-ctxから始まる行です。これを指定する事で、このIPsec通信経路ではLabeled Networkを使用するという事を宣言できます。

spdadd 192.168.11.6 192.168.11.8 any
-ctx 1 1 "system_u:object_r:ipsec_spd_t:s0"
-P out ipsec esp/transport//require;

が、FedoraのInitスクリプトから実行される/etc/sysconfig/network-scripts/ifup-ipsecは、この辺の設定を行うための処理が入っていません。そこで、パッチを当てる事にしました(うげげ…)。

--- ifup-ipsec.orig	2012-12-03 17:41:38.809689669 +0100
+++ ifup-ipsec	2012-12-04 17:52:43.356150231 +0100
@@ -123,6 +123,8 @@ if [ "$ESP_PROTO" = "none" ]; then
     unset SPI_ESP_IN SPI_ESP_OUT KEY_ESP_IN KEY_ESP_OUT SPD_ESP_IN SPD_ESP_OUT
 fi
 
+test -n "$SELINUX_CONTEXT" && SELINUX_CONTEXT="\"${SELINUX_CONTEXT}\""
+
 /sbin/setkey -c >/dev/null 2>&1 << EOF
 ${SPI_AH_OUT:+delete $SRC $DST ah $SPI_AH_OUT;}
 ${SPI_AH_IN:+delete $DST $SRC ah $SPI_AH_IN;}
@@ -164,33 +166,45 @@ EOF
 # This looks weird but if you use both ESP and AH you need to configure them together, not seperately.
 if [ "$ESP_PROTO" != "none" ] && [ "$AH_PROTO" != "none" ]; then
 /sbin/setkey -c >/dev/null 2>&1 << EOF
-spdadd $SPD_SRC $SPD_DST any -P out ipsec
+spdadd $SPD_SRC $SPD_DST any
+            ${SELINUX_CONTEXT:+-ctx 1 1 ${SELINUX_CONTEXT}}
+            -P out ipsec
             ${SPD_ESP_OUT:+esp/$MODE/${TUNNEL_MODE:+$SRC-$DST}/require}
             ${SPD_AH_OUT:+ah/$MODE/${TUNNEL_MODE:+$SRC-$DST}/require}
 	    ;
 
-spdadd $SPD_DST $SPD_SRC any -P in ipsec
+spdadd $SPD_DST $SPD_SRC any
+            ${SELINUX_CONTEXT:+-ctx 1 1 ${SELINUX_CONTEXT}}
+            -P in ipsec
 	    ${SPD_ESP_IN:+esp/$MODE/${TUNNEL_MODE:+$DST-$SRC}/require}
 	    ${SPD_AH_IN:+ah/$MODE/${TUNNEL_MODE:+$DST-$SRC}/require}
 	    ;
 EOF
 elif [ "$AH_PROTO" = "none" ]; then
 /sbin/setkey -c >/dev/null 2>&1 << EOF
-spdadd $SPD_SRC $SPD_DST any -P out ipsec
+spdadd $SPD_SRC $SPD_DST any
+            ${SELINUX_CONTEXT:+-ctx 1 1 ${SELINUX_CONTEXT}}
+            -P out ipsec
             ${SPD_ESP_OUT:+esp/$MODE/${TUNNEL_MODE:+$SRC-$DST}/require}
 	    ;
 
-spdadd $SPD_DST $SPD_SRC any -P in ipsec
+spdadd $SPD_DST $SPD_SRC any
+            ${SELINUX_CONTEXT:+-ctx 1 1 ${SELINUX_CONTEXT}}
+            -P in ipsec
 	    ${SPD_ESP_IN:+esp/$MODE/${TUNNEL_MODE:+$DST-$SRC}/require}
 	    ;
 EOF
 elif [ "$ESP_PROTO" = "none" ]; then
 /sbin/setkey -c >/dev/null 2>&1 << EOF
-spdadd $SPD_SRC $SPD_DST any -P out ipsec
+spdadd $SPD_SRC $SPD_DST any
+            ${SELINUX_CONTEXT:+-ctx 1 1 ${SELINUX_CONTEXT}}
+            -P out ipsec
             ${SPD_AH_OUT:+ah/$MODE/${TUNNEL_MODE:+$SRC-$DST}/require}
 	    ;
 
-spdadd $SPD_DST $SPD_SRC any -P in ipsec
+spdadd $SPD_DST $SPD_SRC any
+            ${SELINUX_CONTEXT:+-ctx 1 1 ${SELINUX_CONTEXT}}
+            -P in ipsec
 	    ${SPD_AH_IN:+ah/$MODE/${TUNNEL_MODE:+$DST-$SRC}/require}
 	    ;
 EOF

この修正により、TYPE=IPSEC として定義されている /etc/sysconfig/network-scripts/ifcfg-* 設定ファイルに「SELINUX_CONTEXT=...」という行が存在すれば、IPsec通信経路を定義する際に、Labeled Networkの定義も同時に行われる事となります。
ipsec_spd_t』というのは、IPsec通信経路に対して定義されているタイプです。現時点ではこれ以外のタイプは定義されていませんので、”お約束”という事にしておきます。

# cat /etc/sysconfig/network-scripts/ifcfg-ipsec0
DST=192.168.1.43
TYPE=IPSEC
ONBOOT=yes
IKE_METHOD=PSK
# KaiGai added them
SELINUX_CONTEXT=system_u:object_r:ipsec_spd_t:s0

うまく設定できていれば、setkey -DPの出力結果の中に『security context: ...』がどーたらという行が出力されているハズです。

# setkey -DP
    :
192.168.1.42[any] 192.168.1.43[any] 255
        out prio def ipsec
        esp/transport//require
        ah/transport//require
        created: Dec  4 16:45:31 2012  lastused: Dec  5 21:26:24 2012
        lifetime: 0(s) validtime: 0(s)
        security context doi: 1
        security context algorithm: 1
        security context length: 33
        security context: system_u:object_r:ipsec_spd_t:s0
        spid=1 seq=0 pid=3500
        refcnt=2

SE-PostgreSQLのインストール

続いて、SE-PostgreSQLのインストールに移ります。
インストール手順は公式ドキュメントにも記載されていますので、Fedoraのパッケージを前提にざっくりと紹介します。

まず、initdbを実行します。

# postgresql-setup initdb
Initializing database ... OK

次に、postgresql.conf を編集して shared_preload_libraries に '$libdir/sepgsql' を追加します。

# vi /var/lib/pgsql/data/postgresql.conf
    :
    :
#max_files_per_process = 1000           # min 25
                                        # (change requires restart)
shared_preload_libraries = '$libdir/sepgsql'

# - Cost-Based Vacuum Delay -

PostgreSQLをシングルユーザモードで起動し、各データベースに初期セキュリティラベルを割り当てます。

# su - postgres
$ for dbname in template0 template1 postgres
  do
    /usr/bin/postgres --single $dbname < \
      /usr/share/pgsql/contrib/sepgsql.sql > /dev/null
  done

これでSE-PostgreSQLの初期設定は完了です。

最後に、TCP/IP経由での接続を受け付けるための設定を行ってサーバを起動します。

# vi /var/lib/pgsql/data/pg_hba.conf
(以下の一行を追加; 設定簡略化のため、とりあえずtrust)
host    all    all    192.168.1.42/32    trust
vi /usr/lib/systemd/system/postgresql.service
(以下の一行を編集し、pg_ctlから-iオプションを渡すようにする)
ExecStart=/usr/bin/pg_ctl start -D ${PGDATA} -s -o "-p ${PGPORT}" -w -t 300
  ↓
ExecStart=/usr/bin/pg_ctl start -D ${PGDATA} -s -o "-i -p ${PGPORT}" -w -t 300

サーバを立ち上げます。

# service postgresql start

試してみる

長かった・・・。

そして、クライアント側からTCP/IP接続でPostgreSQLサーバに接続してみます。

[kaigai@kg1 ~]$ id -Z
unconfined_u:unconfined_r:unconfined_t:s0-s0:c0.c1023

[kaigai@kg1 ~]$ psql -h 192.168.1.43 -U postgres postgres
psql (9.1.6)
Type "help" for help.

postgres=# SELECT sepgsql_getcon();
                    sepgsql_getcon
-------------------------------------------------------
 unconfined_u:unconfined_r:unconfined_t:s0-s0:c0.c1023
(1 row)

ふむふむ、ちゃんとセキュリティコンテキストが切り替わっているように見えます。

それでは、クライアント側で権限を切り替えてからPostgreSQLに接続してみましょう。

[kaigai@kg1 ~]$ runcon -l s0-s0:c0,c2 bash
[kaigai@kg1 ~]$ id -Z
unconfined_u:unconfined_r:unconfined_t:s0-s0:c0,c2
[kaigai@kg1 ~]$ psql -h 192.168.1.43 -U postgres postgres
psql (9.1.6)
Type "help" for help.

postgres=# SELECT sepgsql_getcon();
                   sepgsql_getcon
----------------------------------------------------
 unconfined_u:unconfined_r:unconfined_t:s0-s0:c0,c2
(1 row)

キタコレ!

クライアント側のOS上で変更したプロセスのセキュリティコンテキストが、PostgreSQL側に透過的に伝えられている事が分かります。

なお、以下のようにgetpeercon(3)に失敗している場合は、TCP/IPレベルでの接続には成功したものの、Labeled Networkが機能していません。設定を見直してみてください。

$ psql -h 192.168.0.43 -U postgres postgres
psql: FATAL:  SELinux: unable to get peer label: Protocol not available

まとめ

今回紹介したLabeled Networkの機能を使う事で、SE-PostgreSQLが実現する”OSとDBのアクセス制御の統合”を複数ノードから成る環境にも拡大する事ができるようになりました。
ただ、これは不特定多数からの接続に対して、接続相手のセキュリティコンテキストを受け入れるという設定ではなく、あくまでも、信頼済みのホストが提示するセキュリティコンテキストを受け入れるという形になります。したがって、”誰か分からないけど世界中からアクセスが来る”という場合に使える方法ではなく、自サイトで管理下にあるWebサーバやAPサーバからPostgreSQLサーバに接続する際にセキュリティコンテキストを伝達するための手法という事になります。

現時点では、OS標準のスクリプトがLabeled Networkに対応しておらずスクリプトを修正する必要があったり、初回接続時に(racoonの鍵交換に時間がかかって?)タイムアウトになるといった現象が見られました。
この辺は改良の余地ありなので、引き続き調査してみようと思います。

mod_selinuxのアーキテクチャを考える(後編)

前回の記事では、現在の mod_selinux モジュールの問題点を解決するために複数のタスクキューにワーカースレッドを紐付ける方法を考え、そこから一歩考察を進めて、FastCGIを使ってマルチプロセス実装にすればよりセキュアなWebアプリケーション環境を実現できるのではないか?というアイデアを紹介した。

実際にやってみた。

もちろんSELinuxの権限に紐付けてリクエストの送出先を切り替えるFastCGIモジュールは存在しないので、既存のものに少し手を加える事にする。今回は(たぶん最も広く使われているであろう)mod_fcgidをベースに修正を加える事にする。

mod_fcgidモジュールではFcgidWrapperディレクティブを使って、例えば .php 拡張子を持つURLが指定された場合、FastCGIプロトコルを介して/usr/bin/php-cgiにリクエストを処理させるという指定ができる。
例えばこんな感じ。

AddHandler      fcgid-script     .php
FcgidWrapper    /usr/bin/php-cgi .php

で、mod_fcgidのソースを読んでみて気が付いたのだが、この人はFastCGI常駐プログラムが既に起動しているかどうかをチェックするために、内部テーブルをスキャンして、動作中のプロセスを起動したときのコマンドライン文字列(要はFcgidWrapperで指定した内容)と、これから起動するFastCGI常駐プログラムのコマンドライン文字列を比較するという処理を行っている。
そうすると、仮にラッパープログラムのコマンドライン文字列の一部を、HTTPリクエストを処理するセキュリティコンテキスト(あるいはユーザIDでも可)で置換するような仕組みを作ってあげれば、FcgidWrapperを次のように指定するだけで、後はmod_fcgidが善きに計らってくれる。

FcgidWrapper  "runcon %(SELINUX_CONTEXT) /usr/bin/php-cgi" .php

例えば、HTTPユーザ alice の場合に %(SELINUX_CONTEXT)が "system_u:system_r:httpd_t:s0:c0" に置換されたら、この設定内容は次のコマンドと等価である。おそらく、認証モジュールで設定するようなHTTP環境変数か何かを使うのが汎用的な実装だろう。兎に角、%(SELINUX_CONTEXT)の部分を動的に置換するというのがポイントである。

FcgidWrapper  "runcon system_u:system_r:httpd_t:s0:c0 /usr/bin/php-cgi" .php

もう一点。PHPのような動的コンテンツであれば既存のFastCGIの設定と同様にできるが、静的コンテンツの場合、apache/httpdはhandlerフックの一番最後で他のモジュールが処理できないタイプのHTTPリクエストを処理する…という流れになる。
つまり、ap_hook_handler() に APR_HOOK_LAST を付けて関数ポインタを登録し、静的コンテンツを読むだけのFastCGI常駐プログラムを実行するような拡張が必要になる。

この2点の魔改造を加えた mod_fcgid がこちら。
https://github.com/kaigai/mod_fcgid

KaiGai版 mod_fcgid のインストールと設定

mod_fcgid(魔改造版)のインストールは以下の通り。

$ git clone git://github.com/kaigai/mod_fcgid.git
$ cd mod_fcgid
$ git checkout --track origin/defhnd
$ ./configure.apxs && make
$ sudo make install

Apache/httpdの設定ファイルを編集。ここでは、.php ファイルに対してrunconを介して/usr/bin/php-cgiを実行するのと同時に、静的コンテンツに対してはfcgi-cat(後述)コマンドを実行する。
runconの引数に%(AUTHENTICATE_SELINUX_TYPE)と%(AUTHENTICATE_SELINUX_RANGE)が付いているが、これは実行時にHTTP環境変数と置換される。で、この環境変数を設定するのがauthn_dbd_moduleで、AuthDBDUserPWQueryが2個以上のカラムを返すとき、それらの内容はHTTP環境変数として後続のモジュールで参照する事ができる。

LoadModule fcgid_module      modules/mod_fcgid.so
LoadModule dbd_module        modules/mod_dbd.so
LoadModule authn_dbd_module  modules/mod_authn_dbd.so

DBDriver    pgsql
DBDParams   "host=localhost dbname=web user=apache"

AddHandler      fcgid-script  .php
FcgidWrapper    "/bin/runcon -t %(AUTHENTICATE_SELINUX_TYPE)  \
                             -l %(AUTHENTICATE_SELINUX_RANGE) \
                     /usr/bin/php-cgi" .php
FcgidDocumentWrapper  "/bin/runcon -t %(AUTHENTICATE_SELINUX_TYPE)  \
                                   -l %(AUTHENTICATE_SELINUX_RANGE) \
                           /usr/local/bin/fcgi-cat"

<Directory "/var/www/html">
# BASIC authentication
AuthType    Basic
AuthName    "Secret Zone"
AuthBasicProvider    dbd
AuthDBDUserPWQuery    "SELECT upasswd, selinux_type, selinux_range \
                       FROM uaccount WHERE uname = %s"
Require     valid-user
</Directory>

次に、認証DBのセットアップ。uaccountテーブルにBASIC認証の情報をインプットしておく。$PGDATAはPostgreSQLのDBクラスタのトップディレクトリ。まぁ、COPYコマンドで読めればドコでもいいんですが。

$ htpasswd -c $PGDATA/uaccount.master alice
New password:
Re-type new password:
Adding password for user alice
$ htpasswd $PGDATA/uaccount.master bob
New password:
Re-type new password:
Adding password for user bob
$ createdb web
$ pgsql web
web=# CREATE USER apache;
CREATE ROLE
web=# CREATE TABLE uaccount (
     uname   text primary key,
     upasswd text,
     selinux_type text,
     selinux_range text
 );
CREATE TABLE
web=# GRANT SELECT on uaccount TO public;
GRANT
web=# COPY uaccount(uname,upasswd) FROM '/opt/pgsql/uaccount.master';
web=# COPY uaccount(uname,upasswd) FROM '/opt/pgsql/uaccount.master' DELIMITER ':';
COPY 2
web=# UPDATE uaccount SET selinux_type = 'httpd_t';
UPDATE 2
web=# UPDATE uaccount SET selinux_range = 's0:c0' WHERE uname = 'alice';
UPDATE 1
web=# UPDATE uaccount SET selinux_range = 's0:c1' WHERE uname = 'bob';
UPDATE 1
web=# SELECT * FROM uaccount;
 uname |                upasswd                | selinux_type | selinux_range
-------+---------------------------------------+--------------+---------------
 alice | $apr1$HQtZm8Sz$ZwXm9tx6Kz7g2wiE1LBUZ/ | httpd_t      | s0:c0
 bob   | $apr1$YJWn7hx3$iJyazYClFeFF0n8EgIMSi. | httpd_t      | s0:c1
(2 rows)

次に、静的ドキュメントを参照するFactCGI常駐プログラムのfcgi-catをビルドして /usr/local/bin にインストールする。こやつのビルドには fcgi-devel パッケージも必要です。

$ git clone git://github.com/kaigai/toybox.git
$ cd toybox
$ gcc -g fcgi-cat.c -lfcgi -o fcgi-cat
$ sudo install -m 755 fcgi-cat /usr/local/bin

最後に、SELinuxセキュリティポリシーを追加し、Booleanを調整して設定は終わり。

$ make -f /usr/share/selinux/devel/Makefile webapps.pp
$ sudo semodule -i webapps.pp
$ sudo setsebool httpd_can_network_connect_db on

動かしてみた

お試しとして、以下の内容のPHPスクリプトを実行させてみる。

<?php
echo "<h3>".selinux_getcon()."<h3>\n";
echo "<h3>username = ".$_SERVER["REMOTE_USER"]."</h3>\n";

phpinfo();
?>

ユーザ alice としてアクセスした場合。確かにDBで指定した通り、セキュリティコンテキストの末尾が ":s0:c0" となっている。

一方、ユーザ bob としてアクセスした場合。今度はセキュリティコンテキストの末尾が ":s0:c1" となっている。Server APIの欄が "CGI/FastCGI" となっている点にも注目。

では、静的ドキュメントを参照した場合。2つの画像ファイル test00.jpg と test01.jpg のセキュリティコンテキストはそれぞれ以下のように設定されているため、alice は test00.jpg を参照できるが、test01.jpg は見えないはず。

$ ls -Z /var/www/html
-rw-r--r--. root root unconfined_u:object_r:httpd_sys_content_t:s0 mytest.php
-rwxr--r--. root root unconfined_u:object_r:httpd_sys_content_t:s0:c0 test00.jpg
-rwxr--r--. root root unconfined_u:object_r:httpd_sys_content_t:s0:c1 test01.jpg

これが test00.jpg を参照したケース。正しく読み込めている。

今度は test01.jpg を参照した場合。サーバは 403 Forbidden を返している。ヒャッハー!

画像は省略するが、これをユーザbobで試した場合は逆の結果となる。

もちろん、これらはFastCGIで動作しているので、psコマンドを叩くと常駐プロセスの様子が見える。
確かにユーザ alice と bob それぞれのHTTPリクエストを処理するために、別個のプロセスが起動している。

[root@iwashi ~]# ps -AZ | grep php-cgi
system_u:system_r:httpd_t:s0:c0 23960 ?        00:00:00 php-cgi
system_u:system_r:httpd_t:s0:c1 24062 ?        00:00:00 php-cgi
[root@iwashi ~]# ps -AZ | grep fcgi-cat
system_u:system_r:httpd_t:s0:c0 24098 ?        00:00:00 fcgi-cat
system_u:system_r:httpd_t:s0:c1 24107 ?        00:00:00 fcgi-cat

まとめ

今までの mod_selinux と違い、この方法では全てのDSOモジュールに作用させるという事はできない。
PHPならPHP用の、RubyならRuby用の設定がそれぞれ必要となる。これは確かにデメリット。

その一方で、HTTPユーザ毎にアプリを実行するメモリ空間を分けられるという事は、セキュリティの観点からは非常に大きなアドバンテージである。Web管理者がその気になれば、Apache/httpdでは一切Webアプリの実行を行わず、ある種のアプリケーションサーバとして振る舞うFastCGI常駐プログラムが、HTTPユーザ毎に閉じたコンテナ環境を提供するという事も可能だろう。
セキュリティの強化を目的とするモジュールなので、ここのアドバンテージはでかい、と思う。

mod_selinuxのアーキテクチャを考える(前編)

mod_selinuxapache/httpdのプラグインで、利用者のHTTP要求の認証(BASICとかDIGESTとか)結果に基づいて、PHPスクリプトや静的コンテンツを参照するContents Handlerの実行権限を切り替える。
やっている事は単純で、ap_hook_handlerの先頭で mod_selinux が処理を乗っ取り、一時的にワーカースレッドを作成して権限を縮退させる。で、以降のContents Handlerの実行はワーカースレッドに託されるという形になる。

だが、このデザインだと問題になる事が2点。

  • ap_invoke_handler()が再入可能ではない
  • Bounds domainの範囲でしかドメイン遷移ができない

前者の問題について。CGIが自サーバ内のURLを含むLocation: ヘッダを返した場合など、internal redirectionという処理が走り、cgi_handlerのコンテキストでap_invoke_handler()が再び呼ばれる事がある。
この場合、再びHTTP認証が実行され、その結果、cgi_hanclerのコンテキストとは異なるsecurity contextが指定される事があり得る。が、既にsecurity contextを変更する権限を放棄しているハズなので、詰む。この場合はエラーを返してごめんなさいするしかない。
この問題は、SELinuxに限った話ではなく、CAP_SETUIDを放棄するシナリオでも同じ。

後者の問題について。これはSELinux固有の話で、マルチスレッドのアプリがドメイン遷移を行う場合、遷移先のドメインは bounds domain 制約を満足する必要がある。
(いや、これは元々私の提案だが・・・。)
ドメイン httpd_t が user_webapp_t の bounds domain であるという状態は、user_webapp_t の持っている権限は全て httpd_t も持っているという事。つまり、ドメイン遷移元のhttpd_tは、システム管理系のツール等もWeb経由で実行させると考えると、相当に広範な権限を付与する必要がある・・・事になる。
実際にセキュリティポリシーを書くとなると、これが地味に効いてくる。
端的に言って、書きにくい・・・。(´ω`;;)


まず前者の問題を解決するため、mod_selinuxの構造を上の模式図のように変更してみた。このデザインではsecurity context毎にタスクキューを持ち、ap_handler_hookの先頭でHTTP認証に応じてキューにタスクを振り分けるという形になる。
キューは on demand で生成され、適合する security context が存在しない場合にはTask Launcher Threadにキュー生成要求を行ってこれを作成してもらう。
このデザインだと、ap_invoke_handler()の延長でinternal redirectionが走った場合でも、再帰して呼び出されたap_handler_hookの先頭で行う処理は『適切なキューに処理要求を投げる』だけなので問題は生じない。

・・・と、ここまで実装してハタと気が付いた。

これって本質的には (1)プロトコル処理サーバ と (2)アプリケーションサーバ の分離と同じことだよねと。
だとすると、上のデザインでは各キューに紐付いていたワーカースレッドをワーカープロセスとして分離すると、最初に挙げた2つの問題のうち後者の Bounds domain の制限を取り払う事ができるようになる。

つまり、利用者のcredentials(uid/gid/security context)毎にワーカープロセスをon demandで起動し、HTTPリクエストの解析はapache/httpd自身で行う一方で、静的コンテンツの参照も含めたHTTPリクエストの処理をワーカープロセスで実行するようにすれば、問題は全て解決するのではなかろうか。

幸いな事に、この手のプロトコルとしてFastCGIが既に定義されており、PHPRubyJavaなどメジャーどころの言語はこれに対応している。

...とすると、mod_selinuxを拡張するよりも、mod_fcgidに対する追加機能としてSELinux(やその他のセキュリティ機構を包含する)仕組みを提案した方が、メンテナンスが楽になるのかな、と考え中。

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

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) などで呼びかけてもらえれば。