読者です 読者をやめる 読者になる 読者になる

同期DMAと非同期DMA

GPU PostgreSQL SSD

おっとっと、やらかしてしまった(但し、良い方に)。

PG-Strom + NVMe-Stromでのパフォーマンス計測の際に、SSDからロードしたデータ以外に、例えばテーブル定義情報や定数パラメータといったSQLの実行に必要な情報は一般的なRAM-to-GPU DMAで転送していたのだけども、ココがうっかり同期DMAになっていたために、本来の性能を発揮できないでいた。

そこで、きちんと非同期DMAを実行できるようにコードを修正し、改めてPG-Strom + NVMe-Stromの実行速度を測り直した数字が以下の通り。じゃん。

f:id:kaigai:20160908214519p:plain

ワークロードは変わらず、以下の三種類のクエリを64GB/7億件のテーブルに対して実行した。

  • Q1: 比較的シンプルな検索条件を持つスキャン
  • Q2: 比較的複雑な検索条件を持つスキャン
  • Q3: 文字列マッチ(LIKE句)を持つスキャン

応答時間が概ね42~43secの範囲に収まっているのがお分かりいただけると思う。
これをスループットに直すと、64GB / 43sec = 1524MB/sec のレートでデータを処理できており、Intel SSD 750のカタログスペック 2200MB/s からすると概ね70%程度となる。

応答時間に殆ど差が見られないという事で、これは、GPUで実行するクエリの評価はI/Oよりも短時間で完了するために、非同期DMA転送の裏側に隠ぺいされてしまったと考える事ができる。

CUDAでは非同期で実行する個々のタスク(例えば、RAM=>GPUへのデータ転送、GPU Kernelの実行、など)の順序関係を制御するために、ストリーム(CUstream)という制御構造を持っている。
ある種当然ではあるが、ホスト側から送出されてくるデータを用いて計算しようというGPU Kernelは、少なくともデータ転送が終わらなければ実行できないし、計算結果をデバイス側⇒ホスト側に転送する時も、GPU Kernelの実行が終わっていないと、計算途中のGPU RAMのイメージを送りつけられても困惑する限りである。

CUDAの持つRAM=>GPUのデータ転送用APIには二種類あり、一つは cuMemcpyHtoD() などの同期API。これは、関数が呼び出された時点で呼び出し元をブロックし、RAM-to-GPU DMAを用いてデータを転送する。関数から戻った時点で、既にGPU側のバッファにデータの転送は終わっており、ストリームとは無関係に使用できる。
もう一つは cuMemcpyHtoDAsync() などの非同期API。これは、ストリームにDMA要求が突っ込まれた順番に、非同期にデータ転送を行う。GPU Kernelの実行開始なども、ストリームに要求が突っ込まれた順序を元にした依存関係を壊さないように、ただし、DMAチャネルなどのリソースが空けば待っているタスクはどんどん実行される。

ただし、cuMemcpyHtoDAsync()を用いても必ずしも非同期DMAになるかと言えば、少々落とし穴があり、CPU側のRAMを『ここはDMAバッファだからスワップアウトしないでね』とCUDAランタイムに登録した領域以外を転送元/転送先に指定した場合、黙って同期DMAモードで動作するのである。
今回の場合がそれで、本来はDMAバッファを cuMemHostRegister()で登録しておかねばならないところを忘れていた。Orz

結果、本来であれば次々とデータ転送を行えるところが、ストリームに突っ込んだ cuMemcpyHtoDAsync() が実行可能になるまでブロックされた挙句、同期DMAを行ったものだから、トータルの処理時間が随分と間延びした形になっていた。あーあ。

まぁ、スコアとしては前に計測した時から50%ほど伸びて、対PostgreSQLで見てみると、3~4倍程度のスループットを発揮しているので善しとしよう。

(EN) GpuScan + SSD-to-GPU Direct DMA

GPU SSD PostgreSQL

An article for none-Japanese readers....

What I'm recently working on is a feature to load data blocks from NVMe-SSD to GPU using peer-to-peer DMA.
It allows to bypass the CPU/RAM under a series of data loading process, thus, also allows to reduce number of expensive data copy.
Once data blocks are loaded onto the GPU's RAM, we can process individual records within the blocks, by thousands processor cores, according to the GPU binary which shall be generated based on SQL query.

Here is my longstanding project; PG-Strom that is an extension of PostgreSQL to off-load multiple SQL workloads on GPU devices, transparently. It has been developed for four years, and now supports simple scan, tables join, aggregation and projection.
Its prime focus is CPU intensive workloads, on the other hands, it didn't touch storage subsystem of PostgreSQL because the earlier version of PG-Strom assumes all the data set shall be pre-loaded onto physical memory. No need to say, we had a problem when we want to process a data-set larger than physical RAM.

One my solution is a feature enhancement to support data loading using SSD-to-GPU Direct DMA.
Please assume a use case of a large table scan when its WHERE clause filters out 90% of rows. Usually, DBMS once loads entire data blocks from the storage to RAM, then evaluate individual records according to the WHERE clause. However, this simple but massive calculation is a job GPU can process much more efficiently, and can eliminate 90% of the data prior to loading onto CPU/RAM. It also enables to keep much more valuable data rather than cache out by junk data.

f:id:kaigai:20160825002243p:plain

Once we could load data blocks to GPU prior to CPU, we already have a feature to run tables join and (partial) aggregation on GPU. From the standpoint of CPU, it looks like the storage system gets an intelligence to run a part of SQL workloads, because row-filtering, tables join and aggregation are already done when data stream is arrived at CPU in respond to its i/o requests.

A few days before, the initial version of SSD-to-GPU Direct DMA driver gets working. So, I tried to measure its throughput to load data blocks from a test file to GPU RAM. The result was remarkable.
The "NVMe-Strom" is a record of SSD-to-GPU Direct DMA. Its throughput to load data blocks onto GPU device exceeds about 2200MB/s; which is catalog spec of Intel SSD 750(400GB).
It is also twice faster than a pair of usual read(2) and cuMemcpyHtoDAsync; that it a CUDA API to kick asynchronous RAM-to-GPU DMA. When i/o size is smaller, its system-call invocation cost drives down the throughput furthermore.
f:id:kaigai:20160905215358p:plain

  • CPU: Intel Xeon E5-2670 v3 (12C, 2.3GHz)
  • GPU: NVIDIA Tesla K20c (2496C, 706MHz)
  • RAM: 64GB
  • OS: CentOS 7 (kernel: 3.10.0-327.18.2.el7.x86_64)
  • FS: Ext4 (bs=4096)

In addition, I tried to measure the performance of a few SQL queries:

  1. A large table scan with simple WHERE clause
  2. A large table scan with complicated WHERE clause
  3. A large table scan with LIKE clause

The tables size is about 64GB, contains 700million rows. The result is quite interesting.

f:id:kaigai:20160908214519p:plain
*1

The result of "PG-Strom" (red) shows the performance of GPU acceleration based on the existing storage system of PostgreSQL. It implies the throughput is dominated by i/o, thus unavailable to complete tables scan less than 140sec.
It is valuable when WHERE clause takes complicated expression, but less advantage in other cases.

On the other hands, "PG-Strom + NVMe-Strom" (blue) broke this limitation. In case of simple query, it scans the 64GB table by 42.67sec, thus, its processing throughput is 1535MB/s.
We can expect PG-Strom is more valuable for more CPU intensive workloads, like JOIN or GROUP BY, rather than simple tables scan. I expect the pair of PG-Strom and NVMe-Strom will accelerate this kind of workloads going beyond the existing i/o boundary.
I'm really looking forward to benchmarks of the workloads that contains JOIN and GROUP BY when I could add support of SSD-to-GPU Direct DMA on these features!

QUERY for table construction)

CREATE TABLE t_64g (id int not null,
                    x float not null,
                    y float not null,
                    z float not null,
                    memo text);
INSERT INTO t_64g (SELECT x, random()*1000, random()*1000,
                             random()*1000, md5(x::text)
                     FROM generate_series(1,700000000) x);

postgres=# \d+
                         List of relations
 Schema |       Name       | Type  | Owner  |  Size   | Description
--------+------------------+-------+--------+---------+-------------
 public | t                | table | kaigai | 965 MB  |
 public | t_64g            | table | kaigai | 66 GB   |

QUERY1)
SELECT * FROM t_64g WHERE x BETWEEN y-2 AND y+2;

QUERY2)
SELECT * FROM t_64g WHERE sqrt((x-200)^2 + (y-300)^2+ (z-400)^2) < 10;

QUERY3)
SELECT * FROM t_64g WHERE memo LIKE '%abcd%';

*1:Performance was measured again because asynchronous DMA mode was not enabled in the previous result. Thus, synchronous copy blocked unrelated tasks.

GpuScan + SSD-to-GPU Direct DMA

SSD GPU PostgreSQL

前回の動いた!SSD-to-GPU Direct DMA - KaiGaiの俺メモの記事では、Intel SSD 750とNVIDIA Quadro K1200を使って、Raw-I/OでのSSD-to-GPU Direct DMAが動くところまでを紹介した。

この時点で測定できたSSD-to-GPU Direct DMAのスループットは概ね1400MB/s程度で、VFS経由のスループット1100MB/sを約20%程度上回るものであった。

ただ、Quadro K1200のスペックは以下のようなもので、お世辞にもハイパフォーマンスを期待して使用するタイプのカードではない。*1

チップ GM107GL (CC5.0)
CUDAコア数 512コア, 1000MHz
メモリ容量 4GB GDDR5
メモリ帯域 80GB/s, 128bit

という事で、同じくGPUDirect RDMAに対応した Tesla K20c を使って同じように測定を行ってみた。その結果が以下の通り。

驚いたことに、GPUへのデータロードがIntel SSD 750のカタログスペック通り、2200MB/sまで出ている。
一方で、VFS経由のデータロードは自宅のPCで実行した時と大差ない。

これは当然といえば当然で、①SSD->(VFS)->RAM ②RAM->GPUと二段階のコピーを行わねばならないが、①のSSD->(VFS)->RAMはカタログスペックでも高々2.2GB/sの帯域しか無いにも関わらず、②のRAM->GPUは10GB/sを越える帯域を持っているため、支配項は①の処理であり、ここはGPUの種類に関係のない箇所である。

なお参考までに、VFS経由でSSD->RAMへとデータロードを行った場合のスループットをddコマンドを用いて計測してみた。

基本的にはI/Oサイズが大きいほうが、システムコール呼び出しのオーバーヘッドが少ない分、性能上は有利である。(実際にBlock-I/Oが発行されるのはもっと細かい単位になるが)
ただ、PostgreSQLのブロックサイズは基本8KB、ビルド時のコンフィグによっては32KBまで拡大可能だが、この程度のI/Oサイズでは高速ストレージの応答時間に対するシステムコール呼び出しオーバーヘッドが無視できないレベルになってしまい、スループットが700MB/s程度まで落ち込んでしまっている。
また、もしかしてPage Cacheにコピーを持つ事がオーバーヘッドの主因なのではと思い、O_DIRECT付きのI/Oを試してみたが、こちらは外れ。I/Oサイズを32MB取ったにも関わらず、却ってパフォーマンスの低下を招いてしまった。

で、ここまでが前フリでRaw-I/OでのSSD-to-GPU Direct DMAのおさらい。

昨日、ようやくGpuScanとSSD-to-GPU Direct DMAを統合することができ、全体的にはまだまだバグバグ不安定*2ではあるものの、なんとか一発動かすことができた。

以下のグラフは、64GB/7億件のレコードを持つテーブルに対して3種類のスキャンクエリを実行し、その応答時間を計測したもの。
比較対象は以下の3つで、結果はクエリの応答時間であるので、これが短いほどベター。

各クエリは以下の通り

  1. 比較的シンプルな条件句を持つスキャンクエリ
    • SELECT * FROM t_64g WHERE x BETWEEN y-2 AND y+2;
  2. 比較的複雑な条件句を持つスキャンクエリ
    • SELECT * FROM t_64g WHERE sqrt((x-200)^2 + (y-300)^2+ (z-400)^2) < 10;
  3. 文字列パターンマッチを含むスキャンクエリ
    • SELECT * FROM t_64g WHERE memo LIKE '%abcd%';

まず気になるのが、検索条件の複雑さに関わらず、従来のPG-Stromの応答時間が概ね140s前後に収まっていること。
PG-StromのGPU側処理は非同期で実行されるため、I/Oが十分に遅いようであれば検索条件の評価は完全にI/Oの裏側に埋もれてしまい、そこが処理性能の差としては見えなくなってしまう。
つまり、Intel SSD750 + Ext4上に構築したデータベースのスキャンにおいては、64GB / 140sec = 468MB/s 程度のスループットPostgreSQLのストレージ層の全力であるように見える。*3

一方で、PG-Strom + NVMe-Stromのケース。
比較的条件句がシンプルでDMAバッファが次々に空いたと考えられるQ1のケースでは、64GB / 67.19sec = 975.38MB/s のスループットを発揮している。

これだけでも既にPostgreSQLのシーケンシャルリードの限界を越えているのだが、PostgreSQLのストレージ層がRaw-I/Oの帯域587.38MB/sに対しておよそ20%減の468MB/sで処理できている事を考えると、まだ伸ばせる余地はあるように思える。
データの転送自体はDMAで非同期処理になるとはいえ、このレベルのI/O要求を出し続けるとなると、コマンドの発行だけでも相当なCPU負荷になる。例えば、PostgreSQL 9.6の新機能であるCPU並列にGpuScanが対応し、複数のワーカが同時にSSD-to-GPU Direct DMAを発行し続ければ、NVMe-SSDの理論帯域まで使い切る事も可能かもしれない。

現時点では、スキャンを担当するGpuScanへの対応を行ったところだが、GpuJoinやGpuPreAggのロジックでも全く同じ原理を用いてSSD-to-GPU Direct DMA機能を実装することができる。

元々、CPU負荷が高い方が、PostgreSQLとの差が付きやすいというのはQ2の結果を見ても明らかである。
所詮、GpuScanは条件句の評価にすぎないが、よりCPUインテンシブなJOINやGROUP BYの処理をSSD-to-GPU Direct DMAベースで実装し、パフォーマンスを測定するのが楽しみになってきた。

*1:このカードを購入した経緯は エルザ・ジャパン様の対応が神レベルだった件 - KaiGaiの俺メモを参照。

*2:例外処理で色々とアカン事が起こったりする。:-)

*3:この辺、何かチューニングの余地があるかもしれないが、今回はそこまで掘り下げられていない。

動いた!SSD-to-GPU Direct DMA

SSD GPU


ここしばらく、NVMe-SSDからGPUへとPeer-to-Peer DMAを行うためのLinux kernelドライバを書いている。

これは昨年末のPGconf.JPのLTでアイデアを先に発表したもので、従来は、例えばテーブルスキャンに際して90%の行がフィルタリングされる場合であっても、データをストレージからRAMにロードしていた。しかし、どうせフィルタリングするのであれば、バッファのために利用したRAMのうち90%は無駄である。

基本的なアイデアは、ストレージからのデータロードに際して、CPU側のRAMではなく、GPU側のRAMへロードし、そこで数百~数千コアの計算能力を使って行のフィルタリングや、あるいは、テーブル同士のJOINや集約演算を行ってしまう。そして、これらの前処理が終わった段階でCPU側へデータを書き戻してやれば、CPUから見ると『ストレージからデータを読出したら、既にJOINもGROUP BYも終わっていた』という状態を作り出す事ができる。

まだPostgreSQL側での対応が済んでいないが、やっとドライバだけは動くようになったので簡単な性能テストを行ってみた。

まず、GPUの情報を確認してみる。
搭載しているのはQuadro K1200で、これはドライバの必要とするGPU Direct DMA機能に対応している。

で、デバイスを nvidia-smi -q で確認すると、つらつらとデバイス情報が表示される。
このモデルはGPU RAMを4GB搭載しているが、実はGPU Direct DMAでは搭載RAMの全ての領域を同時に使用できる訳ではない。
ここでBAR1 Memory Usageと書かれたエリアが256MB存在するが、これがPCIバスを通じてホスト側の物理アドレス空間にマップする事のできる大きさで、いわば4GBのGPU RAMを256MB分の大きさの窓から覗くようなものである。

[kaigai@magro ~]$ nvidia-smi -q

==============NVSMI LOG==============

Timestamp                           : Thu Aug 25 00:49:35 2016
Driver Version                      : 352.93

Attached GPUs                       : 1
GPU 0000:01:00.0
    Product Name                    : Quadro K1200
    Product Brand                   : Quadro
           :
    FB Memory Usage
        Total                       : 4095 MiB
        Used                        : 9 MiB
        Free                        : 4086 MiB
    BAR1 Memory Usage
        Total                       : 256 MiB
        Used                        : 1 MiB
        Free                        : 255 MiB
    Compute Mode                    : Default
           :

なお、このBAR1領域は、Tesla K40/K80やM40といった上位モデルでは、物理アドレスサイズを越える16GBが用意されており、これらのGPUを持っているリッチメンであればBAR1領域の小ささに悩む必要もないハズである。

とりあえず、今は手元にQuadro K1200しかないので、これでトライ。

データ転送のパターンとしては、32MBのバッファを6個確保し、ファイルの内容を順にこれらのバッファに非同期転送を行う。非同期転送が完了したタイミングでCUDAからコールバックが呼び出されるため、バッファが空いたらすかさず次の非同期転送をキューイングしていく。

本来であれば、GPUで何か処理を実行してデータサイズを減らしてからCPU+RAMへ書き戻すのが正しい姿であるが、まだGPU側で処理すべき事を実装できていないので、ここは単純にSSDから読み出した内容を同じバッファへと書き戻す事とした。

通常のVFS経由でのデータロードだと、上記の図のようになる。
read(2)でSSDからバッファにデータをロードし、これをcuMemcpyHtoDAsync()とcuMemcpyDtoHAsync()を使ってCPU RAM⇔GPU RAM間でPCI-Eバスを介して移動させる。

一方、SSD-to-GPU Direct DMAを使用した場合、データの流れはシンプルになる。
SSDからGPUPCI-Eバスを介してPeer-to-Peerでデータ転送が行われ、その後、GPU RAM⇒CPU RAMへの転送が行われる。

かなり乱暴に考えても、CPU RAMに一旦コピーする手間が省ける分は処理性能が稼げるはずである。

これを、OSのキャッシュの影響を避けるため、16GB搭載のマシン上で40GB、100GBの二つのファイルに対して実行してみた。

結果は以下の通り。

性能指標としては、データ転送量を一連の処理時間で割ったスループットを使用している。
SSD-to-GPU Directのケースが最も早く、概ね1.4GB/s程度のスループットを出している。ただし、Intel SSD 750 (400MB)のカタログスペックはシーケンシャルリード 2200MB/s と記載してあるので、これに比べるとまだデバイスの性能を発揮しているとは言い難い。
とりあえず、ドライバが動きましたという段階なのでアレだが、今後、どこで引っかかっているかは調べる必要がある。

一方、VFS経由でデータをロードした場合は1.1GB/s程度のスループットに留まっている。これは、SSD⇒CPU RAMへのデータ転送が余分に必要となっている分、不可避の差とも言える。ただ、これはPostgreSQLでのワークロードを反映しているとは言えない。
というのも、1.1GB/sを記録したケースというのは、read(2)にバッファサイズである32MBを与えた場合であって、システムコール呼び出しに関わるオーバーヘッドが極めて少ないと言えるため。
通常、PostgreSQLのI/Oの単位は8KBで、コンフィグで設定を変えたとしても32KBまでしか増やす事ができない。
こうなると、40GBや100GBといった大きさのファイルをスキャンする際、特にNVMe-SSDのような高速デバイスを使う場合だと、システムコールの呼び出しが律速要因となってしまうため、大幅にスループットが低下しているのが判る。

現在、PostgreSQL側の対応も実装中であるので、単純なI/Oだけでなく、実際にPostgreSQL+PG-Stromにクエリを食わせてみた時の性能値についても追って報告したい。

オレオレ Demand Paging

PostgreSQL GPU

現在の PG-Strom のアーキテクチャは、PostgreSQLの各バックグラウンドプロセスが個別にCUDAコンテキストを作成し、GPUバイスメモリを作るという構成になっている。
これは、設計の単純化、特にエラーパスのシンプル化により、全体的なソフトウェアの品質が低い時には開発効率の観点からは意味のあるデザインではあるが、一方で、以下のような問題も同時に抱えている。

  • CUDAコンテキストの初期化には200ms~400ms程度のオーバーヘッドを要するため、数秒で終わる程度のクエリ処理ではこのコストは無視できない。
  • 一方、CUDAコンテキストの初期化・破棄の回数を抑えるために、一度構築したCUDAコンテキストをキャッシュしておくと無駄にGPU RAMを保持し続ける。CUDAコンテキストあたり~90MB程度のGPU RAMを消費する*1ので、同時並行セッション数が増えてくるとワーキングメモリが取れなくなる。
  • GPUリソースの過剰割当て。あるバックエンドが巨大バッチの実行中にGPU RAMを使い過ぎ、その後、別のセッションが新たに始まっても、巨大バッチ実行中のバックエンドからGPU RAMを奪う事ができない*2

f:id:kaigai:20160619082709p:plain

なので、ある意味先祖還りなのだが、上の図のようにGPUやCUDAに関わる部分だけを別プロセスに分割して、GPUリソースの管理や同時実行数のコントロールを行えるようにしたい。特に、使わなくなったGPUリソースを直ちに解放できるようにする事と、CUDAコンテキスト作成のオーバーヘッドを軽減する事は、次のPostgreSQLでCPU+GPUハイブリッド並列を実装するにあたって必須とも言える作業である。

とはいえ、PostgreSQLバックエンド ⇔ CUDAサーバ間のデータの受け渡しという新たな問題が出てくる。

  • PostgreSQL起動時に獲得する静的共有メモリ領域はサイズが固定であるので、サイジングが難しい。少なすぎれば途中で out of memory だし、大きすぎれば他に使うべきメモリ領域を圧迫する。
  • 動的に共有メモリを獲得した場合、複数のバックエンド間でアドレスを共有できない。したがって、一般的に、データの受け渡しには共有メモリセグメント先頭からのオフセット表現を用いるか、mmap(2)MAP_FIXEDを使って固定アドレスでセグメントがマッピングされるよう強制するしかない。
  • しかしmmap(2)MAP_FIXEDを使用した場合でも、ポインタを受け取った先でその共有メモリセグメントが既にマップされているかどうかは毎回バリデーションが必要で、これならオフセット表現と大差ない。

....という所で悶々としていたのだが、この辺の問題を解決する方法を考えてみた。

  • mmap(2)MAP_FIXEDを使って共有メモリセグメントをマップするアドレスは固定にする。
  • 各共有メモリセグメントは各々1GBなどの固定長。各セグメントの状態だけは静的共有メモリ上に保持。
  • SIGSEGVSIGBUSシグナルハンドラを使用して、これらの共有メモリセグメントへの参照が必要になった時点でmmap(2)する。
  • 実際にmmap(2)するまでは、使用する可能性のある仮想アドレス空間にはPROT_NONEでダミーの領域をマップしておく。

....なーんでもっと早く気付かなかったかなぁ?という位単純な話だが、図を使って説明すると以下のようになる。


f:id:kaigai:20160619082718p:plain

初期状態。PostgreSQLの各バックエンドはpostmasterと呼ばれる親プロセスからfork(2)するが、その時点で、オレンジ色の静的共有メモリと、PROT_NONEでmmap(2)した仮想アドレス空間を持つよう初期化されている。
各共有メモリセグメントは、segment_idrevisionを持ち、これによって共有メモリセグメントの存在・不存在の状態を管理する。mmap_ptrはこれをマップすべき領域の仮想アドレスである。
で、各バックエンドはfork(2)して作られるので、この仮想アドレス空間を引き継ぐことになる。

f:id:kaigai:20160619082759p:plain

次に、あるバックエンドが共有メモリセグメントを作成し、それを自分の仮想アドレス空間にマップする。
共有メモリアロケータは、ここにより小さなメモリブロック(chunk)を割り当てる事ができるが、この時点では共有メモリセグメントはまだ他のバックエンドにmmap(2)されていないため、このポインタは不可視である。

f:id:kaigai:20160619082809p:plain

では、このポインタを他のバックエンドに渡した時に、どのような振る舞いを見せるか。
他のバックエンドでこのポインタの参照はSIGSEGVを引き起こす。デフォルトではプロセスのクラッシュを引き起こすが、シグナルハンドラを追加してやる事で当該ポインタを含む共有メモリセグメントをmmap(2)する事ができる。

f:id:kaigai:20160619082829p:plain

そうすると、ポインタを受け取ったバックグラウンドの側でもその領域が可視となる。
ポインタで参照されていた領域があたかも初めからそこに存在したかのように振る舞う事ができるので、『共有メモリセグメントを参照する可能性のあるポインタを毎回バリデーション…』といった、面倒でバグの温床となりがちなプログラムを書かずに済む。

まず手始めに、この仕組みをCUDAプログラムのキャッシュに適用してみた。
このコードは、PG-Stromが自動生成したCUDAプログラムをビルドして生成されるバイナリイメージを共有メモリ上に残しておき、次に同じクエリを実行する際のコンパイルを省略するためのもの。
CUDAプログラムのコンパイルはバックグラウンドワーカーにより非同期で行われるため、共有メモリ上に置かれている

postgres=# EXPLAIN ANALYZE SELECT count(*) FROM t0 NATURAL JOIN t1;
dmaBufferAttachSegmentOnDemand: pid=24048 got Segmentation fault,
then attached shared memory segment (id=0 at 0x7de0d6124000, rev=1)

                               QUERY PLAN
------------------------------------------------------------------------------
 Aggregate  (cost=3102815.03..3102815.04 rows=1 width=0)
            (actual time=10407.955..10407.955 rows=1 loops=1)
   ->  Custom Scan (GpuPreAgg)  (cost=14139.24..2872442.40 rows=256 width=4)
                                (actual time=4249.929..10407.925 rows=3 loops=1)
         Reduction: NoGroup
         ->  Custom Scan (GpuJoin) on t0  (cost=10139.24..2852814.83 rows=100000080 width=0)
                                          (actual time=71.376..9900.973 rows=100000000 loops=1)
               GPU Projection:
               Outer Scan: t0 (actual time=11.080..5052.385 rows=100000000 loops=1)
               Depth 1: GpuHashJoin, HashKeys: (t0.aid)
                        JoinQuals: (t0.aid = t1.aid)
                        Nrows (in:100000000 out:100000000, 100.00% planned 100.00%)
                        KDS-Hash (size: 9.16MB planned 13.47MB, nbatches: 1 planned 1)
               Inner Buffer: (13.47MB), DMA nums: 150, size: 2020.76MB
               ->  Seq Scan on t1  (cost=0.00..1935.00 rows=100000 width=4)
                                   (actual time=0.019..13.689 rows=100000 loops=1)
 Planning time: 14.645 ms
 Execution time: 10774.832 ms
(14 rows)

以下のようにデバッグメッセージが出ており、SIGSEGVをシグナルハンドラで受け取り、
結果、共有メモリセグメントをオンデマンドでマッピングした事を示している。

dmaBufferAttachSegmentOnDemand: pid=24048 got Segmentation fault,
then attached shared memory segment (id=0 at 0x7de0d6124000, rev=1)

めでたしめでたし。

*1:Tesla K20cで観察してみた結果。Quadro K1200では15MB程度だったので、デバイスによって変わるのかも。

*2:自発的にGPU RAMを解放できれば良いが、PG-Strom関連の処理を行った後CPU-onlyなタスクが30分回ってる、という場合、クエリ全体の完了を待つ以外には誰も使っていないGPU RAMを解放する手立てがない

PL/CUDAとmatrix型

GPU PostgreSQL

PG-Stromには↓のような利点があるが、

  1. SQLから自動的にGPUバイナリ命令列を生成するため、GPUプログラミングを意識する必要がない、
  2. PostgreSQLの行指向データを用いるため、既存DBからデータの移行を必要としない。

その裏返しとして、同時に↓のような特徴も持っている。

  1. NULL値チェックや四則演算ごとのオーバーフローチェックなどSQLに由来する動作のため、専用に設計されたGPUプログラムほど高速には動作できない。
  2. 行指向であるためデータの密度が低く、目的のデータにアクセスするために何度もRAMアクセスを繰り返す必要がある。

比較対象がSQLであれば、これらの欠点には目をつぶっても、数百~数千コアの計算能力で殴りつけて十分なアクセラレーション性能を得る事ができる。(というか、GPUにデータを供給する足回りの方が先に悲鳴を上げる)

が、例えばRなどのように、専用の統計解析パッケージに対して計算能力で優位に立つ事を考えると、計算の中核部分などはSQL由来のアレやコレやは省く事で、GPU本来のパフォーマンスを引き出してやりたいところである。

そこで実装してみたのが、PL/CUDA言語とmatrix型。

PL/CUDAはSQL関数を実装するための言語で、CREATE FUNCTION構文で定義したSQL関数の定義部分をそのままCUDAのkernel関数としてビルド。関数の引数を自動的にGPUへ送出してユーザ定義の処理ロジックを並列実行し、実行結果をまたCPU側へ書き戻す。

引数にはfloatやtextなどが利用できるが、新たにmatrix型というデータ型を定義した。これの実体は単精度浮動小数点型の2次元配列であるが、全ての値が非NULLである事を保証しているので、行列の(i,j)要素をダイレクトにアクセスできる。要は、計算密度の高い部分には密度の高いデータを供給するようにしようという話である。

f:id:kaigai:20160531211753p:plain

試しにCREATE FUNCTIONでPL/CUDA関数を実装してみる。簡単な例という事で、行列同士の乗算を行うGPU関数を実装してみた。

CREATE OR REPLACE FUNCTION matrix_gpu_mul(matrix, matrix)
RETURNS matrix
AS $$
#plcuda_prep
#plcuda_num_threads 1
  if (!pg_matrix_sanitychecks(kcxt, arg1) ||
      !pg_matrix_sanitychecks(kcxt, arg2))
    PLCUDA_ERROR_RETURN(StromError_InvalidValue);

  MatrixType *X = (MatrixType *)arg1.value;
  MatrixType *Y = (MatrixType *)arg2.value;
  MatrixType *R;
  size_t      len;

  if (X->width != Y->height)
    PLCUDA_ERROR_RETURN(StromError_InvalidValue);

  len = sizeof(cl_float) * (size_t)X->height * (size_t)Y->width;
  if (kplcuda->results_bufsz < len)
    PLCUDA_ERROR_RETURN(StromError_DataStoreNoSpace);

  retval->isnull = false;
  retval->value = (varlena *)results;
  R = (MatrixType *) results;
  pg_matrix_init_fields(R, X->height, Y->width);

#plcuda_begin
#plcuda_num_threads matrix_gpu_mul_num_threads
  MatrixType *X = (MatrixType *)arg1.value;
  MatrixType *Y = (MatrixType *)arg2.value;
  MatrixType *R = (MatrixType *)retval->value;
  cl_uint     x_height = X->height;
  cl_uint     y_height = Y->height;
  cl_float   *xval;
  cl_float   *yval;
  cl_uint     i, j, k, nloops;

  assert(X->width == Y->height);
  nloops = X->width;
  i = get_global_id() % x_height;
  j = get_global_id() / x_height;
  xval = X->values + i;
  yval = Y->values + j * y_height;

  if (get_global_id() < x_height * y_height)
  {
    cl_float sum = 0.0;

    for (k=0; k < nloops; k++)
    {
      sum += (*xval) * (*yval);
      xval += x_height;
      yval ++;
    }
    R->values[get_global_id()] = sum;
  }

#plcuda_end
#plcuda_results_bufsz matrix_gpu_mul_results_bufsz
$$ LANGUAGE plcuda;

何ヶ所か#plcuda_で始まる見慣れないディレクティブがあるが、これがPL/CUDAのディレクティブで、#plcuda_prepから始まるコードブロック、#plcuda_beginから始まるコードブロックがそれぞれ、GPU kernel関数の実体として差し込まれる事になる。
基本的に、#plcuda_prepブロックはバッファの初期化を、#plcuda_beginアルゴリズムの中核部分の処理を、そしてここでは使っていないが#plcuda_postブロックは結果バッファへの書き戻しと、3ステップでGPU kernel関数を実行する事を想定している。
ただ、Dynamic Parallelismを使って他のGPU kernel関数をバンバン起動する事もできるので、別にこの3ステップに拘る必要はない。

コードブロックの中で指定されている#plcuda_num_threadsは、GPU kernel関数を起動する時のスレッド数を指定する。バッファ初期化を行う#plcuda_prepブロックは1スレッドで起動され、行列積の計算を行う#plcuda_begin...#plcuda_endブロックは、別のSQL関数であるmatrix_gpu_mul_num_threadsの返り値と同じ数のスレッドで起動される。
行列積の大きさは入力された行列の大きさによって変わるので、予め特定の値で決め打ちする事ができないためである。

末尾で指定している#plcuda_results_bufszは、結果バッファの大きさを指定する。これは定数でもよいが、行列積のように出力結果が入力行列の大きさによって変わる場合には、別のSQL関数を呼びだして結果バッファを動的に決められるようにすべきである。

で、これを実行した結果、以下のような計算結果が得られる。
パッと見た感じ、単純な3x3行列同士の演算であるが、裏方ではGPU用のコードをビルドし、引数の行列を2つGPUへ転送して計算を行い、結果をCPU側へ書き戻している。

postgres=# select matrix_gpu_mul('{{2,0,0},{0,2,0},{0,0,2}}'::matrix,
                                '{{1,2,3},{4,5,6},{7,8,9}}'::matrix);
         matrix_gpu_mul
--------------------------------
 {{2,4,6},{8,10,12},{14,16,18}}
(1 row)

plcuda_function_source関数を使えば、ここで定義したPL/CUDA関数がどのようなGPU kernel関数に置き換えられ、ビルドされているのかを確認する事ができる。
以下のコード中、/* ---- code by pl/cuda function ---- */で挟まれた部分が、PL/CUDA関数の定義から差し込まれた部分である。

postgres=# select plcuda_function_source('matrix_gpu_mul'::regproc);
                       plcuda_function_source
---------------------------------------------------------------------
 #include "cuda_common.h"                                           +
 #include "cuda_matrix.h"                                           +
 #include "cuda_plcuda.h"                                           +
                                                                    +
 STATIC_INLINE(void)                                                +
 __plcuda_matrix_gpu_mul_prep(kern_plcuda *kplcuda,                 +
               void *workbuf,                                       +
               void *results,                                       +
               kern_context *kcxt)                                  +
 {                                                                  +
   pg_matrix_t *retval __attribute__ ((unused));                    +
   pg_matrix_t arg1 __attribute__((unused));                        +
   pg_matrix_t arg2 __attribute__((unused));                        +
   assert(sizeof(*retval) <= sizeof(kplcuda->__retval));            +
   retval = (pg_matrix_t *)kplcuda->__retval;                       +
   assert(retval->isnull || (void *)retval->value == results);      +
   arg1 = pg_matrix_param(kcxt,0);                                  +
   arg2 = pg_matrix_param(kcxt,1);                                  +
                                                                    +
   /* ---- code by pl/cuda function ---- */                         +
   if (!pg_matrix_sanitychecks(kcxt, arg1) ||                       +
       !pg_matrix_sanitychecks(kcxt, arg2))                         +
     PLCUDA_ERROR_RETURN(StromError_InvalidValue);                  +
   MatrixType *X = (MatrixType *)arg1.value;                        +
   MatrixType *Y = (MatrixType *)arg2.value;                        +
   MatrixType *R;                                                   +
   size_t      len;                                                 +
   if (X->width != Y->height)                                       +
     PLCUDA_ERROR_RETURN(StromError_InvalidValue);                  +
   len = sizeof(cl_float) * (size_t)X->height * (size_t)Y->width;   +
   if (kplcuda->results_bufsz < len)                                +
     PLCUDA_ERROR_RETURN(StromError_DataStoreNoSpace);              +
   retval->isnull = false;                                          +
   retval->value = (varlena *)results;                              +
   R = (MatrixType *) results;                                      +
   pg_matrix_init_fields(R, X->height, Y->width);                   +
   /* ---- code by pl/cuda function ---- */                         +
 }                                                                  +
                                                                    +
 KERNEL_FUNCTION(void)                                              +
 plcuda_matrix_gpu_mul_prep(kern_plcuda *kplcuda,                   +
             void *workbuf,                                         +
             void *results)                                         +
 {                                                                  +
   kern_parambuf *kparams = KERN_PLCUDA_PARAMBUF(kplcuda);          +
   kern_context kcxt;                                               +
                                                                    +
   assert(kplcuda->nargs == kparams->nparams);                      +
   INIT_KERNEL_CONTEXT(&kcxt,plcuda_prep_kernel,kparams);           +
   __plcuda_matrix_gpu_mul_prep(kplcuda, workbuf, results, &kcxt);  +
   kern_writeback_error_status(&kplcuda->kerror_prep, kcxt.e);      +
 }                                                                  +
                                                                    +
 STATIC_INLINE(void)                                                +
 __plcuda_matrix_gpu_mul_main(kern_plcuda *kplcuda,                 +
               void *workbuf,                                       +
               void *results,                                       +
               kern_context *kcxt)                                  +
 {                                                                  +
   pg_matrix_t *retval __attribute__ ((unused));                    +
   pg_matrix_t arg1 __attribute__((unused));                        +
   pg_matrix_t arg2 __attribute__((unused));                        +
   assert(sizeof(*retval) <= sizeof(kplcuda->__retval));            +
   retval = (pg_matrix_t *)kplcuda->__retval;                       +
   assert(retval->isnull || (void *)retval->value == results);      +
   arg1 = pg_matrix_param(kcxt,0);                                  +
   arg2 = pg_matrix_param(kcxt,1);                                  +
                                                                    +
   /* ---- code by pl/cuda function ---- */                         +
   MatrixType *X = (MatrixType *)arg1.value;                        +
   MatrixType *Y = (MatrixType *)arg2.value;                        +
   MatrixType *R = (MatrixType *)retval->value;                     +
   cl_uint     x_height = X->height;                                +
   cl_uint     y_height = Y->height;                                +
   cl_float   *xval;                                                +
   cl_float   *yval;                                                +
   cl_uint     i, j, k, nloops;                                     +
   assert(X->width == Y->height);                                   +
   nloops = X->width;                                               +
   i = get_global_id() % x_height;                                  +
   j = get_global_id() / x_height;                                  +
   xval = X->values + i;                                            +
   yval = Y->values + j * y_height;                                 +
   if (get_global_id() < x_height * y_height)                       +
   {                                                                +
     cl_float sum = 0.0;                                            +
     for (k=0; k < nloops; k++)                                     +
     {                                                              +
       sum += (*xval) * (*yval);                                    +
       xval += x_height;                                            +
       yval ++;                                                     +
     }                                                              +
     R->values[get_global_id()] = sum;                              +
   }                                                                +
   /* ---- code by pl/cuda function ---- */                         +
 }                                                                  +
                                                                    +
 KERNEL_FUNCTION(void)                                              +
 plcuda_matrix_gpu_mul_main(kern_plcuda *kplcuda,                   +
             void *workbuf,                                         +
             void *results)                                         +
 {                                                                  +
   kern_parambuf *kparams = KERN_PLCUDA_PARAMBUF(kplcuda);          +
   kern_context kcxt;                                               +
                                                                    +
   assert(kplcuda->nargs == kparams->nparams);                      +
   if (kplcuda->kerror_prep.errcode != StromError_Success)          +
     kcxt.e = kplcuda->kerror_prep;                                 +
   else                                                             +
   {                                                                +
     INIT_KERNEL_CONTEXT(&kcxt,plcuda_main_kernel,kparams);         +
     __plcuda_matrix_gpu_mul_main(kplcuda, workbuf, results, &kcxt);+
   }                                                                +
   kern_writeback_error_status(&kplcuda->kerror_main, kcxt.e);      +
 }                                                                  +
                                                                    +

(1 row)


今回は、まずPL/CUDA関数を定義し、実行できるようになったところまで。
次はもう少し実際的なケースで、計算量の多いものをPL/CUDA関数で実装する例をご紹介したい。

エルザ・ジャパン様の対応が神レベルだった件

GPU 日常

雑文です。

現在取り組んでいる SSD-to-GPU ダイレクト機能の実装には、PostgreSQL/PG-Strom側の機能拡張だけれなく、NVMe SSDからGPU RAMへのDMAを実行する Linux kernel ドライバの開発が必要になる。

Linux kernelにはDMAを実行するためのインフラが既に多数揃っているので、ドライバの開発自体はそれほど大仕事ではないのだが、GPUがその機能に対応している必要がある。

NVIDIA提供のドキュメントによると、
GPUDirect RDMA :: CUDA Toolkit Documentation

GPUDirect RDMA is available on both Tesla and Quadro GPUs.

と、あり、いわゆるコンシューマ向け廉価製品であるGTXでは対応していない。

対応していないというのは、GPU上のRAMをホストアドレス空間にマップするためのAPIである nvidia_p2p_get_pages() がエラーを返してしまうので、それ以上は如何ともし難いという事である。

int nvidia_p2p_get_pages(uint64_t p2p_token,
                         uint32_t va_space_token,
                         uint64_t virtual_address,
                         uint64_t length,
                         struct nvidia_p2p_page_table **page_table,
                         void (*free_callback)(void *data),
                         void *data);

試しに、手元で利用可能なGPU何種類かでトライしてみたところ

  • × GTX 750Ti
  • × GTX 980
  • ○ Tesla K20c

という結果。GTX980は割とハイエンドのモデルではあるのだが、それでも対応してないモノは対応していない。

会社で使う分には Tesla K20c があるので良いのだが、問題は週末プログラマの開発環境。
特に今年のゴールデンウィークは長いので、"動作確認・デバッグは連休明けまでお預け" なんて事になると精神衛生上も大変によろしくない。

ので、上記の GPU Direct 機能に対応していて、かつ、できるだけ廉価な製品を買う事にした。

Teslaシリーズはさすがに高くて手が出ないので、ワークステーション向け Quadro のラインナップから選択。エントリーモデルかミドルレンジ程度ならなんとか手が出る。

NVIDIA Quadro シリーズ | カテゴリー製品情報 | 株式会社 エルザ ジャパン

とはいえ、今回購入する事になった Quadro K1200 は4万円程度の品で、個人で買うには覚悟の要る品物。『TeslaとQuadroで対応って書いてあるけど、あれってハイエンドモデルだけだから(テヘペロ』みたいのが一番困るので、日本でのNV社代理店であるエルザ・ジャパン様に問い合わせてみた。

質問)
Quadro K1200 の購入を検討しているのですが、GPUDirect RDMAには対応しているでしょうか?

NVIDIA社のドキュメントには以下の記述があります。
http://docs.nvidia.com/cuda/gpudirect-rdma/index.html

GPUDirect RDMA is available on both Tesla and Quadro GPUs.

回答)
GPU Directにつきましては、サポートされていますので、下記のアドレスを参照してください。
http://www.nvidia.com/object/compare-quadro-gpus.html

ただ、製品ごとの機能マトリックスには『GPUDirect for Video』という項目が掲載されており、しかもK1200にはチェックが付いていないという内容であったので、再度確認してみた。

当方が気にしているのは、以下の GPU Direct 機能の対応可否です。
https://developer.nvidia.com/gpudirect

Using GPUDirect, multiple GPUs, third party network adapters,solid-state drives (SSDs) and other devices can directly read and write CUDA host and device memory, eliminating unnecessary memory copies

とあるように、SSDGPUへの直接データ転送を行う Linux kernel ドライバの開発に使用したいと思っているのですが、『本当にこれを買ってしまっていいの?』 という疑問が腹落ちしておりません。

お手数ですが、再度確認をお願いできないでしょうか?

その後、

失礼しました。

実際に動作確認をしますので、少しお時間をいただけますでしょうか?
後日、ご連絡いたします。

という連絡があり、その数時間後に、

別の部署にて確認しましたところ、動作はするようですが、
K1200 の IB 4ノードクラスタより、1ノードで、Tesla の方が速いそうです。

と、実際に動作確認を行った上で回答を頂いた。

GPUは開発・デバッグ用に使いたいだけなので、性能は全く気にしていない。なので、早速ポチっと注文。日曜日の夕方には届いたので早速PCに装着してみた。

f:id:kaigai:20160417192855j:plain
これまで使っていたGTX 750Tiを取り外し、Quadro K1200を装着する。

f:id:kaigai:20160417193139j:plain
Quadro K1200とIntel 750 SSDを同じPCI-Eバス上に装着。
これで SSD-to-GPU ダイレクトを使用する前提条件が整った・・・ハズ。

試しに、作成途中のカーネルモジュールをロードし、テストプログラムで ioctl(2) を叩いてみる。

$ sudo insmod nvme-strom.ko
$ ./driver_test -p 4 /opt/nvme/testfile
vaddr=0x701720000 length=4194304
ioctl(STROM_IOCTL_PIN_GPU_MEMORY) = 0

OK、0が返ってきたという事は、Quadro K1200上で nvidia_p2p_get_pages() が正しく動作している。

$ dmesg | tail -80
[ 6933.359672] nvme-strom: P2P GPU Memory (handle=18446612149830398528) was mapped
  version=65537, page_size=1, entries=64
[ 6933.359675] nvme-strom:   H:0000000701720000 <--> D:00000000e0180000
[ 6933.359676] nvme-strom:   H:0000000701730000 <--> D:00000000e0190000
[ 6933.359677] nvme-strom:   H:0000000701740000 <--> D:00000000e01a0000
[ 6933.359677] nvme-strom:   H:0000000701750000 <--> D:00000000e01b0000
                   :

ログメッセージにも、デバイスの仮想アドレスと物理アドレスの対応が表示されており、GPUのページテーブルを正しく取得できたという事を示唆している。これで勝つる。

という訳で、こんな質問を投げるマニアは一体何人おるんだというニッチな問い合わせに対して、しかも小売価格で高々4万円の製品を売るために、わざわざ動作検証まで行っていただいたエルザ・ジャパン様のユーザ対応に大変感動した次第である。