オレオレ Demand Paging

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

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

なので、ある意味先祖還りなのだが、上の図のように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などの固定長。各セグメントの状態だけは静的共有メモリ上に保持。
• SIGSEGVやSIGBUSシグナルハンドラを使用して、これらの共有メモリセグメントへの参照が必要になった時点でmmap(2)する。
• 実際にmmap(2)するまでは、使用する可能性のある仮想アドレス空間にはPROT_NONEでダミーの領域をマップしておく。

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

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

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

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

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

まず手始めに、この仕組みを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型

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)要素をダイレクトにアクセスできる。要は、計算密度の高い部分には密度の高いデータを供給するようにしようという話である。

試しに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関数で実装する例をご紹介したい。

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

雑文です。

現在取り組んでいる 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

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

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

その後、

失礼しました。

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

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

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

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

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

これまで使っていたGTX 750Tiを取り外し、Quadro K1200を装着する。

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のページテーブルを正しく取得できたという事を示唆している。これで勝つる。

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

GTCで喋りました

という訳で、GTCで喋ってきました。

20160407_GTC2016_PgSQL_In_Place from Kohei KaiGai

www.slideshare.net

今までの発表とは少し趣を変えて、PG-Stromそのものの説明よりも、現実世界のワークロードを実行するときにどういった使い方があり得るか、どういった効果が得られるかに主眼を置いた内容。

伝統的にDBMSは、正に字面の通り、データベースをマネジメントするソフトウェアである訳なので、インデックスを張って目的のレコードを高速に取り出したり、テーブル同士を結合する事は得意だが、（一応SQLの構文で書く事はできるとはいえ）大量の計算をこなすには向いていない。
少なくとも、データ解析専用に設計されたツール（例えばR言語）を使って計算させた方が高速で、しかも統計解析系のパッケージも揃っているので、普通はデータベースからデータをエクスポートした上でデータ解析処理を行う。

では、PG-Stromを入れることで計算処理が早くなれば、この前提が変わるか？というのが今回のお題。

■ DB内で解析処理を行う事のメリット

• サーバ～クライアント間でデータを移動せずにすむ。数GB単位の大きさになると無視できない。
• 通常はクライアント側よりも強力なサーバハードウェアの計算能力を使用できる。
• 常に最新のデータセットに対して解析処理を行うことができる。

■ DB内で解析処理を行うことのデメリット

• サーバ側が十分な計算能力を持っている必要がある。
• DBMSがデータ解析処理に対応するよう設計されている必要がある。
• 専用APと比べると、ライブラリ類の整備が不十分。

ライブラリ類の充実はともかくとして、計算能力に関してはいかほどのものか。

今回の実験のターゲットは、創薬に関連したワークロードの一つで、大量の化合物の中からできるだけ特性の異なった化合物を選択するというもの。これにより、実際に化合物を合成して特性を調べるなど、"お金のかかる" ステップを少なくしたいというモチベーションがある。

使用したアルゴリズムはMAX-MIN法と呼ばれるもの。実際には何個かのクエリを繰り返し実行するのだが、ワークロードの中核、最も計算が重い部分は赤点線で囲った距離計算の部分。
化学物質の特性を数値パラメータ化し、比較対象との"距離"を計算する。
基本的にはこの距離が大きな化学物質を順に選択することで、ある一群からピックアップした k 個の化学物質の特性ができるだけ他と異なるものになる。

これを、SQLによるデータエクスポート+R言語でのローカル計算、SQLで実行+R言語で結果参照 with/without PG-Strom の３パターンでそれぞれトライしてみた。
結果は、PG-Stromなしのパターンが圧倒的に遅く、R言語のローカル実行と、SQL+PG-Stromのリモート実行が概ね同じという結果になった。

そのため、現状でも『わざわざデータをEXPORTしてローカルで計算しなくても、SQLで全部計算させて結果だけ取ってきた方がデータ管理が楽ですよ』というのは言えるだろう。

一方で、他の並行セッションでは、よりHPCに近く計算能力を必要とする人工知能の研究成果が発表されており、『SQLを高速に処理する』がためにGPUを使用するPG-Stromの設計では専用にチューンしたアプリケーションに太刀打ちすることはできない。
（する必要があるのかどうか、という議論は当然あるが）

しかし、データに最も近い場所でCUDAプログラムを実行しうるプラットフォーム機能を既に有しているというアドバンテージを活かすのであれば、OLTP/OLAPという伝統的DBの世界から外れたところのワークロードについて考えてみる価値はあるだろう。おそらく、必要な要素技術は既に揃っている。

検討してみたいのは、前回のエントリでも書いたアイデアで、CUDAのロジックをSQLで記述するための機能。

PostgreSQLでは、SQL関数を記述する言語を拡張モジュールによって定義する事ができる。
したがって、CUDAによって記述されたSQL関数を実現するのはそれほど難しい話ではない。

SQL上で hogehoge(matrix, vector) などと定義された関数であれば、実行時には、第一引数に行列をコピーしたバッファのポインタを、第二引数にベクトルをコピーしたバッファのポインタを渡すようにすればよい。フラットな単純配列に対してCUDAカーネルを起動すれば、あとは普通のHPC屋さんがやっている内容になる。

現状のSQL->CUDAへのコード変換の場合、SQLに由来するNULLチェックや四則演算の都度発生するオーバーフローチェックが入っているため、計算処理に重きを置くのであれば相当に効率の悪いコードになっている。

何から何までCUDAのコードを書かせるのはアレだが、R言語などでもパッケージ化されていて比較的良く使われる関数、例えば kmeans() 相当がCUDAのフルスピードで実行できるようになったらどうだろうか？

SQLでkmeans法のアルゴリズムを記述した場合、繰り返し処理の記述性が高くなく無駄な処理が入っていることもあり、現状では PG-Strom を用いてもRのkmeans()関数の1.5倍程度の時間を要している。（PG-Strom無しだと実行完了を諦めるレベルだが）

CUDAのフルスピードで行列演算を実行できれば、計算部分の処理速度を現状から更に一桁向上させる程度はできるだろう。
RDBMS的にはニッチもニッチでいい所の機能だが、『明らかに高速化の効果を実感できる』機能として、前倒しトライしてみるのも良いかなと思えてきた。

GTCに来ています

今年もサンノゼで開催されている GPU Technology Conference 2016 に参加しています。

# なお、当方の発表『In-Place Computing on PostgreSQL: SQL as a Shortcut of GPGPU』は木曜日の予定

キーノートでは、NVIDIA社CEOのJen-Hsum Huang氏より"ディープラーニング向けプロセッサ"と銘打って、新アーキテクチャPascalを搭載したTesla P100が発表に。また、CUDA8.0のリリースが6月である事も発表された。
（一方で、事前に噂されていたPascalベースのコンシューマ向けGTXモデルへの言及はなかった）

このPascalアーキテクチャベースのTesla P100であるが、Inside Pascal: NVIDIA’s Newest Computing Platformに記載のある通り、

• 150億トランジスタ
• 3584 CUDA Core
• 16GB RAM (HBM2; 720GB/s)
• FP64 5.3TFlops

...等々のお化けプロセッサである。

プログラムを書く上で、Maxwell世代から変わっている点、新機能が追加になった点がある。

• SM(Streaming Multiprocessor)あたりのCUDA Core数が減っている

128個 -> 64個になった。一方でSM数は56個に増えており、これでトータル3584CUDA Coreという訳だが、SMあたりのレジスタファイル容量(256KB)は据え置き、共有メモリ容量(64KB)は33%増加*1であるので、システム全体として見ると、大幅な増強となっている。

また、SMあたりのコア数が少なくなった一方で、レジスタ、共有メモリの容量が据え置きになっているという事は、ある特定のSMに対して一度に投入できるタスクの数が増え、それにより、例えばDRAMアクセス待ちの間に他のタスクにスケジュールして計算コアの使用率を上げる事ができると思われる。

PG-Stromのようにnone-coalescedなメモリアクセスが多いワークロードでは嬉しいかも。

• 64bit浮動小数点のAtomic演算サポート

PG-Stromには嬉しい新機能。現状、PostgreSQLの数値計算系の関数は全て64bit浮動小数点で、標準偏差や分散、共分散などを集計する時には、内部的には cmpxchg 命令を使った atomic 演算のエミュレーションを行わざるを得ないので、H/Wによりネイティブの atomic 演算がサポートされれば集計系の性能向上が期待できる。

• プリエンプションの対応

仮想化環境での利用を見込んだ機能か？上述のSMに投入できるブロック数が増えたという事も併せて考えると、Pascal世代ではGPUの同時並行利用が大いに促進されるような気がする。

• デマンドページング

従来は、予め必要となるバッファサイズを予測してCUDA Kernelを起動する前にデバイスメモリのアロケーションが必要であったが、これにはある程度マージンを取る必要があるので、実際には無駄になる領域が少なからずあった。GPU側でもPage Fault→デバイスメモリ割当てという流れが可能になる事で、リソースの有効利用が可能となる。
一方で、Page Faultのコストは無視できないほど高いので、cudaMemPreFetchAsync()などCUDA8.0で対応となる新APIにより、プログラムがヒントを与えてやる必要がある。

• HBM2

一目で分かるように、スループットが大幅に引き上げ。現行世代の３倍弱。
ただし、これはメモリアクセスがcoalescedな場合の性能向上で、none-coalescedなメモリアクセスはレイテンシが律速要件になるため、現行世代と比較して10%程度の向上との事。

行指向データ構造を使うPG-Stromとしてはちょっと悲しい所だが、対策としては、例えばDRAMからデータを一旦共有メモリにロード(ここはcoalescedなアクセスで)。その後の共有メモリへのアクセスはL1キャッシュと同じレイテンシなので、そんなに気にはならないはず。

SMあたりのCUDAコア数が減った事で、こういう最適化も考えられる。

さて、会場でSさんとHさんに『PG-Stromの上でMachine Learningのアプリとか動かせるようにならないんですかね～？』と無茶振りをされる。

現状、"PostgreSQLと同じ使用感"で使える事を優先するために、本来は効率の悪い行指向データをリッチな計算リソースで無理やり並列処理しているので、専用のアプリケーションでの処理と比べるとどうしても分が悪い。

ただ『データのある場所で計算を行う』というのは正しい方向性なので、SQLで数式を書いたら自動的にコード生成⇒GPU実行、という流れに拘らなければ、SQLクエリ実行の過程でデータ密度の高いデータ構造を作り出し、これをGPUで実行するMachine Learning用の（でも、何か他の数値計算の）アプリケーションに渡してやるという事はそんなに不自然ではないのかな、とは思う。

例えば、32bit浮動小数点型の2次元配列・NULLなしというデータ構造であればぎっしりデータが詰まっており、先頭からのオフセットだけでアクセスすべきデータを特定できる。つまり、行指向ならではの不都合はない。これを仮に matrix 型と名付けて、

SELECT my_deep_leraning(SELECT make_matrix(x0, x1, x2, ..., xN)
                          FROM data_source
                         WHERE 抽出条件
                         GROUP BY グループ化条件);

みたいなクエリで実行できれば、、、、って事にはなるだろう。

ただ、当然に考えねばならん課題もあり、
・生成するMATRIXがGPUのRAMサイズに収まらない場合
・粗行列の場合、本当に単純な配列表現でいいのか？
・一行に対して複数のCUDA Kernelで処理する事になるので、PG-Stromにとってはwhat's new

面白いアイデアなので、もう少し考えてみたいところではある。

*1:たぶんブロックあたりのサイズの事を言っている。SMあたり容量ならMaxwellの上の方は96KBなので

SSD-to-GPU Peer-to-Peer DMAとバッファ管理（その２）

前回の続き。

PCI-E接続のSSDからのP2P DMAでCPU/RAMを介さずにGPU RAMへデータを転送するという要素技術自体は目新しいものではない。

かつてFusion-io（現: SunDisk）もやっていたし、NVMe規格に準拠したものであれば標準のドライバに少し手を加えてP2P DMAを実行する事も可能である。実際にやっている人もいる。

問題は、PostgreSQLでの利用シーンを考えた時に、常にストレージからのP2P DMAを実行すればよい、というワケではなく、バッファにキャッシュされたデータの扱いや並行プロセスとの相互作用を考慮して、矛盾なく、かつ、性能的にもリーズナブルである方法を考えねばならない。

先ず、大前提として考えねばならないのが、PostgreSQLの追記型アーキテクチャとMVCC(Multi-version concurrency control)の仕組み。
かいつまんで言うと、各レコードには『そのレコードを作成した人のトランザクションID(xmin)』と『そのレコードを削除した人のトランザクションID(xmax)』という値が記録されており、参照側トランザクションIDとの大小関係や、これらのトランザクションがcommitされているか否かといった情報を元に、レコードの可視/不可視を判定する。

これらのルールはトランザクション分離レベル（READ COMMITTEDやREPEATABLE READ）によって異なるが、それなりに複雑なロジックを処理するのでできれば避けたい。（さらに言えば、同時並行で動いているトランザクション状態にも依って判定が変わるため、GPUで処理する事はできない。）

そこで使われるのが visibility map と、all_visible フラグ。

PostgreSQLでは8KB～32KBの大きさのブロック単位でデータを管理しており、ブロックに包含されるレコードが全て、あらゆるトランザクションから明らかに可視である時には all_visible フラグがセットされる。このフラグが立っている時には、個別レコードのMVCCチェックは必要ない。
さらに、これはブロック自体とは別管理の visibility map(VM) と呼ばれるデータ構造と連動しており、VMのビットが立っているブロックは、（ブロック自体を読み込まなくても）all_visibleフラグがセットされている事が分かる。*1

そもそもGPUではall_visible=0であるブロックからレコードを抽出できないので、all_visible=0であるブロックはCPU側でレコードを抽出してRAM->GPUへと転送しなければならない。したがって、この時点でSSD-to-GPU P2P DMAの対象にできるのはall_visible=1のブロックである事が分かる。*2

VMの状態から次に読もうとするブロックがGPUでの処理が可能である事が分かった。次に、このブロックが既にPostgreSQLのバッファに載っているかどうかを確認する。
NVMe SSDは確かに高速なデバイスではあるが、RAMと比較すると文字通り"桁違いの"遅さであるので、同じPCI-Eバスを介してデータ転送を行うのであれば、普通にRAM->SSDへのDMAを行った方が合理的である。

したがって、SSD-to-GPU P2P DMAの対象となるブロックは ①all_visibleフラグが立っている ②まだshared_bufferにロードされていない という条件を満たすものである事が分かる。

さて、並行するプロセスが介在する場合を考える。ある時点で対象となるブロックがshared_bufferに載っていない事が分かり、P2P DMAのリクエストをキューに入れたとする。しかし、それとはお構いなしに、並行プロセスが当該ブロックをOSからロードし、更新処理を行ってしまうかもしれない。
ただし、この事自体は問題とはならない。更新されたレコードは、P2P DMAを行っているトランザクションが開始した時点では未だコミットされておらず、あたかも存在しないものであるかのように扱っても構わないからである。
したがって、OS buffer/Storageに保持されている古いバージョンのブロックをGPUへ転送し、そこで処理を行っても問題とはならない。

注意しなければならないのは、並行プロセスによって更新されたレコードが、P2P DMAの完了までにOSに書き戻される場合。

PostgreSQLのストレージ書き込み（smgrwrite）はBuffered Writeなので、書き出された内容はまずOSのPage Cacheに保持される。これは単なるメモリコピーであるが、Page CacheからGPU RAMへとDMA転送を行っている途中にPage Cacheが更新され、旧バージョンと新バージョンが混じり合った状態のデータがGPUに転送されてしまうと最悪である。

ファイルシステムにもよるが、write(2)システムコールの処理中は inode->i_mutex により排他ロックを取る事になっているようなので、不完全な状態のデータを使用する事を避けるには、P2P DMAの実行開始*3から転送完了までの間は、NVMeドライバの側でも inode->i_mutex を取ってやらないとダメ。

さらに、inode->i_mutex で排他処理を行った場合でも、P2P DMAの実行より前に、並行プロセスによって更新済みページが既にOSバッファに書き出されたという場合が考えられる。このようなケースでは、GPU側でまずブロックの all_visible フラグをチェックし、GPU RAMに転送された時点でこれが all_visible=0 にクリアされていた場合にはいったん諦め、CPU側でMVCCチェックと改めてRAM上に確保したDMAバッファにレコードを積んでおく事を要求する。

そもそも、P2P DMAリクエストから転送完了までのわずかな時間の間に、並行プロセスが当該ブロックをOSから読み出し、更新を行い、それをまたOSバッファに書き戻すという一連の処理を実行する頻度はかなり低いと想定される上に、更新済みブロックは並行プロセスの手で既に shared_buffer にロードされているので、この再実行によって新たにI/Oの負荷が発生する事はない。

気を付けなければならないのは、OSのPage Cacheへの書き出しとDMAの処理がカチ合わないよう、ここだけはきちんと排他処理を踏んでやる事である。

*1:これはIndexOnlyScanでも使われており、PG-Stromとは無関係に元々ある機能

*2:とはいえ、一般的な業務アプリでは90%の行は更新されない（by SAP）らしいので、通常はall_visible=1となるブロックが大半であろう。

*3:リクエストキューに積んだ時点である必要はない

SSD-to-GPU Peer-to-Peer DMAとバッファ管理（その１）

昨年の暮れ、JPUGカンファレンスのLTで『SQL+GPU+SSD=∞』と題したスピーチを行った。

SQL+GPU+SSD=∞ (Japanese) from Kohei KaiGai

www.slideshare.net

これはかいつまんで言えば、ストレージからデータをCPU+RAMへとロードするより前に一旦GPUへとデータを転送し、そこで不要なデータを削ぎ落してからCPU+RAMへと渡してやる事で、CPU負荷の軽減とRAMの有効活用が計れるというアイデアである。
実装としては、PCI-Eデバイス間でのP2P DMA機能を利用する事によってNVMe SSDの特定ブロックからGPU RAM上の特定の領域へDMAを実行するというものなので、ここは別に新しくも何ともない。

以下の図は、従来の仕組みにおけるデータの流れを示したもの。

SSDから読み出されたデータは先ずCPU+RAMにバッファされ、それをGPUに送出してJOIN処理を実行、結果を受け取るという一連の流れである。しかし、①SSD=>RAMにデータが転送される時、本来は検索条件に合致せず削除されるはずのタプルや、問合せの中で参照されないカラムも含まれている。つまり、ゴミデータもRAM上のバッファを消費し、もしかしたらその煽りで別のキャッシュされていたブロックが弾き飛ばされてしまう可能性だってある。

もしPG-StromがSSD-to-GPU P2P DMA機能に対応すればどのように変わるか？

SSDから読み出されたデータは、CPU+RAMへの転送より前に、先ずGPUへと転送される。ここで検索条件や参照列のチェックを行い、先ず、不要な行と列は削除する。場合によっては、SSDから直接読み出したテーブルと結合する事になるInner側のテーブルが既にCPU側から送られており、不要行・列の削除に留まらず、JOINまで行った状態でCPU+RAMへ結果を転送できるかもしれない。
こうすると、CPU+RAMの視点から見た時、『SSDからデータブロックを読みだしたと思ったら、なんとJOINまで既に終わっていたでござる』という状況を作る事ができるようになる。

ただ、ここで考えねばならないのがバッファ管理。
PostgreSQLの全てのブロックの内容がストレージ（NVMe SSD）と常に同期されていれば話は簡単なのだが、そうは問屋が下ろさない。

PostgreSQLはデータブロックの書き出しをbuffered writeで行うので、ストレージの内容よりも新しいデータが、PostgreSQL自身のバッファ（shared buffer）、あるいはOSのバッファ（page cache）に存在する可能性がある。
したがって、バッファされていない、即ちストレージに最新のデータが存在する場合にはSSD-to-GPU P2P DMAを行うとしても、バッファされているデータとの整合を考慮しなければならない。
つまり、GPUを使おうとするプログラムは、バッファリング状態を確認した上で適切なDMA元を指定してやる必要がある。

先ず、目的とするブロックがPostgreSQLのshared buffer配下に存在するか否かの確認。これは BufTableLookup() 関数を用いれば確認する事ができ、既に載っていれば、shared bufferをソースとするDMAを実行する。
既にCPU+RAM側で保持しているデータに対しては、GPUでデータの振るい落としを行おうがそもそもRAM使用量の節減効果はない事と、そもそもRAMはSSDよりも速いデバイスだから、という事である。

PostgreSQLのshared bufferに載っていない場合、Linuxのpage cache状態の確認はユーザプロセスからは行えないので、ドライバ側で適切にハンドリングする必要がある。これも find_get_page() 関数を用いれば確認する事ができ、既に載っていれば、page cacheをソースとするDMAを実行する。
この2パターンはどちらも本体RAMをソースとするDMAで、それに該当しない場合、NVMe SSDの該当ブロックをソースとするDMAを実行する事になる。

ちょっと辛いのが、『特定ファイルの特定オフセット位置』をブロックデバイス上のセクタ番号に変換するためのインターフェースがVFSで規定されていない事。なので、NVMe SSDドライバにSSD-to-GPU対応機能を追加するとしても、利用可能なファイルシステムを何個か絞った上で、ファイルオフセット⇒ブロック番号への変換関数を呼びだして対処する必要がある。*1
Ext4ならext4_get_block、XFSならxfs_get_blocks関数なんてものが定義されているので、これらを呼び出せばそう難しいものではないハズだが…。

もう一つ考慮しなければならないのが、競合する更新処理をどのように捌いていくのか。これはエントリを分けて記す事にしたい。

*1:LVMやSoft-RAIDを使用する場合はもう一段階の変換が必要か