本章では、PL/CUDA言語を用いて、GPUで実行可能なネイティブプログラムをSQL関数として実装する方法について説明します。

PL/CUDA概要

内部的に、PG-StromはSQL構文を元にCUDA言語によるGPUプログラムを生成し、これを実行時コンパイルによってGPU用命令バイナリを生成します。 CUDAとはNVIDIA社の提供するプログラミング環境で、C言語に似た構文を用いてGPUで並列実行可能なプログラムを記述する事ができます。 SQL構文からCUDAプログラムへの変換プロセスは内部的なもので、ユーザの視点からは、どのようなGPU用プラグラムが生成、実行されるのかを意識する必要はありません。

一方、PostgreSQLではCREATE LANGUAGE構文を用いてSQL関数の記述に用いるプログラミング言語を追加する事ができます。 PL/CUDAとはCREATE LANGUAGE構文に対応した言語ハンドラで、SQLを元にPG-Stromが自動生成するGPUプログラムだけでなく、ユーザが実装した任意のGPUプログラムをSQL関数として実行する事が可能となります。

SQL関数の引数には、数値型や文字列型、行列型など、PG-Stromのサポートするデータ型を使用する事ができますが、これらはPL/CUDA実行系が自動的にGPU側へデータを転送するため、データベースとGPU間のデータロードについて意識する必要はありません。また同様に、PL/CUDA関数の戻り値(可変長データ型である場合を含む)もGPU側からCPU側へと書き戻され、SQL関数の戻り値として整形されます。

また、PL/CUDA関数の引数としてgstore_fdwを用いて定義した外部表を使用する事ができます。この場合、データは既にGPUにロード済みであるためPL/CUDA関数呼び出しのたびにデータロードを行う必要はなく、またPostgreSQL可変長データの長さ制限である1GBよりも大きなデータを使用する事ができます。

これらの特徴により、ユーザはGPUやデータベースとの間のデータの入出力といった定型的な処理に煩わされる事なく、統計解析ロジックの実装や高速化といった生産的な作業に注力する事ができます。

PL/CUDA Overview

CREATE FUNCTION構文を用いてPL/CUDA関数を定義すると、この関数の定義部を含むCUDAプログラムのソースコードを作成し、これをターゲットGPU向けにビルドします。 このCUDAプログラムは、引数の受け渡しと結果を返却するための補助的なコードを含む以外は、一般的なCUDAランタイムを用いたソフトウェアと全く同一で、CUDAの提供する各種のライブラリをインクルード/リンクする事も可能です。

PL/CUDA関数を用いて作成したネイティブのCUDAプログラムは、PostgreSQLバックエンドの子プロセスとして実行されます。 したがって、PostgreSQLとは独立したアドレス空間と、OSやGPUのリソースを持つ事になります。 CUDAプログラムには、ホストシステム上で実行されるホストコードと、GPU上で実行されるデバイスコードを含みます。ホストコードはC言語でプログラミング可能なあらゆるロジックを実行可能ですので、セキュリティ上の観点から、PL/CUDA関数の定義はデータベース特権ユーザに限定されています。

以下に単純なPL/CUDA関数の例を示します。 この関数は、同じ長さのread型配列を二つ引数に取り、そのドット積をfloat型で返却します。

CREATE OR REPLACE FUNCTION
gpu_dot_product(real[], real[])
RETURNS float
AS $$
#plcuda_decl
#include "cuda_matrix.h"

KERNEL_FUNCTION_MAXTHREADS(void)
gpu_dot_product(double *p_dot,
                VectorTypeFloat *X,
                VectorTypeFloat *Y)
{
    size_t      index = get_global_id();
    size_t      nitems = X->height;
    float       v[MAXTHREADS_PER_BLOCK];
    float       sum;

    if (index < nitems)
        v[get_local_id()] = X->values[index] * Y->values[index];
    else
        v[get_local_id()] = 0.0;

    sum = pgstromTotalSum(v, MAXTHREADS_PER_BLOCK);
    if (get_local_id() == 0)
        atomicAdd(p_dot, (double)sum);
    __syncthreads();
}
#plcuda_begin
{
    size_t      nitems;
    int         blockSz;
    int         gridSz;
    double     *dot;
    cudaError_t rc;

    if (!VALIDATE_ARRAY_VECTOR_TYPE_STRICT(arg1, PG_FLOAT4OID) ||
        !VALIDATE_ARRAY_VECTOR_TYPE_STRICT(arg2, PG_FLOAT4OID))
        EEXIT("arguments are not vector like array");
    nitems = ARRAY_VECTOR_HEIGHT(arg1);
    if (nitems != ARRAY_VECTOR_HEIGHT(arg2))
        EEXIT("length of arguments mismatch");

    rc = cudaMallocManaged(&dot, sizeof(double));
    if (rc != cudaSuccess)
        CUEXIT(rc, "failed on cudaMallocManaged");
    memset(dot, 0, sizeof(double));

    blockSz = MAXTHREADS_PER_BLOCK;
    gridSz = (nitems + MAXTHREADS_PER_BLOCK - 1) / MAXTHREADS_PER_BLOCK;
    gpu_dot_product<<<gridSz,blockSz>>>(dot,
                                        (VectorTypeFloat *)arg1,
                                        (VectorTypeFloat *)arg2);
    rc = cudaStreamSynchronize(NULL);
    if (rc != cudaSuccess)
        CUEXIT(rc, "failed on cudaStreamSynchronize");

    return *dot;
}
#plcuda_end
$$ LANGUAGE 'plcuda';

PL/CUDA実行系は、#plcuda_begin#plcuda_endで囲まれた部分にSQL関数の引数の受け渡しを行う処理を付加して、CUDAプログラムのエントリポイントを作成します。 #plcuda_decl#plcuda_beginで囲まれた部分は、GPUデバイス関数やその他のホスト関数を宣言するためのブロックで、ソースコード上では上記のエントリポイントより前に配置されます。

CUDAプログラムのエントリポイントでは、arg1arg2、... という形でSQL関数の引数を参照する事ができます。

上記の例では、real[]配列型であるarg1およびarg2がエントリポイントへ渡され、VALIDATE_ARRAY_VECTOR_TYPE_STRICTマクロによってNULLを含まない32bit浮動小数点型の1次元配列であるかどうかを検証しています。

返り値も同様に、SQLデータ型に相当するCUDA Cデータ型をエントリポイントから返します。 エントリポイントがreturnで値を返さない場合(または明示的にexit()で終了コード1を返した場合)、PL/CUDA関数はNULL値を返却したものとして扱われます。

上記のサンプルプログラムでは、SQL関数から受け取ったreal型配列を検証した後、cudaMallocManagedで結果バッファを獲得した後、GPUカーネル関数であるgpu_dot_productを呼出してドット積を計算しています。

この関数の実行結果は以下の通りです。ランダムに生成した10,000個の要素を持つベクトル同士の内積を計算しています。

postgres=# SELECT gpu_dot_product(array_matrix(random()::real),
                                  array_matrix(random()::real))
             FROM generate_series(1,10000);
 gpu_dot_product
------------------
 3.71461999509484
(1 row)

PL/CUDAの構造

PL/CUDAの関数定義は、#plcuda_decl#plcuda_begin、および#plcuda_endの各ディレクティブによって分割されたコードブロックから構成されます。各コードブロックには各々の目的に応じたユーザ定義のCUDA Cコードを記述する事ができ、これらは、PL/CUDA言語ハンドラとの引数及び結果の受け渡しを行うロジックと結合し一個のソースファイルへと再構成されます。

#plcuda_decl
  [...any declarations...]
#plcuda_begin
  [...host code in the entrypoint...]
#plcuda_end

#plcuda_declより始まるコードブロックは、__host__および__device__属性を持つCUDA C関数や変数を完全な形で記述する事ができます。 このコードブロックは、最終的に構築されるソースファイル上で、#plcuda_begin...#plcuda_endブロックを含むエントリポイントよりも前方に位置します。

また、CUDA Cの#include構文を用いて外部ヘッダファイルをインクルードする場合には、このコードブロックに記述するようにして下さい。

#plcuda_beginより始まるコードブロックは、ホストコードであるエントリポイント関数の一部として組み込まれます。したがって、関数名や引数の型などを記述する事はできません。 エントリポイント関数は、当該コードブロックに制御が移る前に、パイプを介してSQL関数の引数をPostgreSQLバックエンドから受信し、arg1arg2、…という名前で参照できるようセットアップを行います。

これらの変数は、SQLデータ型に応じて以下の表に示すCUDA Cとしての表現を持ちます。

SQLデータ型 CUDA Cデータ型 説明
reggstore void * Gstore_fdw外部表のOID
real float 32bit浮動小数点型
float double 64bit浮動小数点型
その他のインライン型 Datum intdateなど
固定長ポインタ型 void * uuidなど
可変長データ型(varlena) varlena * textreal[]など

PL/CUDA言語ハンドラは、上記のコードブロックから一個のCUDA Cソースファイルを作成し、宣言時または実行時に1度だけnvccコンパイラでこれをビルドしCUDAプログラムを生成します。 #plcuda_includeディレクティブを含む場合はCUDA Cソースファイルが実行時にしか確定しないため、実行時にのみこれをビルドします。ただし、同一内容のCUDAプログラムがビルド済みである場合にはこれを再利用します。

PL/CUDA Callflow

SQLからPL/CUDA関数を呼び出すと、PL/CUDA言語ハンドラはビルド済みのCUDAプログラムを起動し、パイプを通じてSQL関数の引数をコピーします。引数はCUDAプログラム内の引数バッファに格納され、これらはarg1arg2などの名前で参照する事が可能です。

可変長データ型などCUDA Cプログラム上でポインタとして表現されるデータ型は、引数バッファへの参照として初期化されます。引数バッファはcudaMallocManaged()によって獲得されたmanaged memory領域であるため、当該ポインタはホスト⇔デバイス間の明示的なDMAなしに使用する事ができます。

引数がreggstore型を持つ場合は特殊です。これは本来Gstore_Fdw外部テーブルのOID(4バイト整数)を表現するデータ型ですが、PL/CUDAの引数として与えられた場合はGstore_fdwが獲得しているGPUデバイスメモリへの参照へと置き換えられます。 引数はGstoreIpcMappingオブジェクトへの参照として初期化され、GstoreIpcMapping::mapにはGstore_Fdw外部テーブルの確保したGPUデバイスメモリをマップしたアドレスが入ります。 当該領域を物理的に保持しているGPUデバイスIDはGstoreIpcHandle::device_idを、当該領域の長さはGstoreIpcHandle::rawsizeを参照してください。

typedef struct
{
    cl_uint     __vl_len;       /* 4B varlena header */
    cl_short    device_id;      /* GPU device where pinning on */
    cl_char     format;         /* one of GSTORE_FDW_FORMAT__* */
    cl_char     __padding__;    /* reserved */
    cl_long     rawsize;        /* length in bytes */
    union {
#ifdef CU_IPC_HANDLE_SIZE
        CUipcMemHandle      d;  /* CUDA driver API */
#endif
#ifdef CUDA_IPC_HANDLE_SIZE
        cudaIpcMemHandle_t  r;  /* CUDA runtime API */
#endif
        char                data[64];
    } ipc_mhandle;
} GstoreIpcHandle;

typedef struct
{
    GstoreIpcHandle h; /* IPChandle of Gstore_Fdw */
    void       *map;    /* mapped device pointer */
} GstoreIpcMapping;

PL/CUDA関数の処理結果を返すには、SQLデータ型に対応するCUDA Cのデータをエントリポイントからreturnで返却します。 値を明示的に返却しない場合、CUDA Cでのデータ型がポインタ型でありNULLを返した場合、あるいはCUDAプログラムがexit(1)によりステータスコード1で終了した場合は、PL/CUDA関数はSQLに対してnullを返したものとして扱われます。

PL/CUDAリファレンス

本節はPL/CUDA関数のディレクティブ、および関連するSQL関数のリファレンスです。

PL/CUDAの得意不得意

PL/CUDA関数が呼び出されると、その背後でCUDAプログラムが起動され、CUDAプログラムはGPUデバイスの初期化を行います。これらの一連の処理は決して軽いものではなく、例えば単純なスカラー値の比較を行うようなロジックをPL/CUDA関数で実装し、10億行のフルテーブルと同時に使用するという使い方は推奨されません。

一方で、ひとたびGPUデバイスの初期化が完了すれば、GPUの持つ数千プロセッサコアを利用して大量データを高速に処理する事が可能です。特に、繰り返し計算により最適パラメータを計算する機械学習や統計解析のように、ワークロードに占める計算の割合が大きな問題に適すると言えるでしょう。

処理すべきデータが増加すると、CUDAプログラムとのデータの受け渡し方法にも注意が必要です。 PostgreSQLは配列型をサポートしており、整数型や実数型のデータを高々数百万個程度受け渡すのであれば手軽な方法です。 しかし、配列型を含むPostgreSQLの可変長データの最大長は1GBであるため、これより巨大なデータの受け渡しにはデータの分割など工夫が必要です。また、SQL関数の引数をセットアップするのはPostgreSQLバックエンドプロセスで、この処理はシングルスレッドで動作するためGB単位のメモリ操作には相応の時間を要します。

データサイズが数百MBを越えてきた段階で、Gstore_Fdw外部テーブルの利用を検討してください。 Gstore_Fdwを通して予めGPUデバイスメモリにデータをロードする事で、PL/CUDA関数の呼び出し時に長大な引数をセットアップする必要がなく、また、GPUデバイスメモリ容量が許す限り、GBを越えるサイズのデータを保持する事が可能です。

PL/CUDAディレクティブ

#plcuda_decl

このディレクティブは__host__および__device__属性を持つCUDA C関数や変数を含むコードブロックを開始します。PL/CUDA言語ハンドラは、CUDAプログラムのソースファイル上で、このコードブロックをエントリポイントよりも前にそのままコピーします。

このディレクティブの使用は任意ですが、エントリポイントから呼び出すべきGPUカーネル関数を宣言しなければPL/CUDA関数を使用する意味はありませんので、通常は一つ以上のGPUカーネル関数を含む事になります。

#plcuda_begin

このディレクティブは、CUDAプログラムのエントリポイントを構成するコードブロックを開始します。 CUDAプログラムは、受け取ったPL/CUDA関数の引数をarg1arg2、...という変数名で参照可能となるよう初期化を行った上で、コードブロックへと制御を移します。当該コードブロックはホストコードであり、CPUで動作する制御ロジックや、GPUカーネルを呼出しての計算処理を記述する事ができます。

結果を返すには、CUDA Cのreturn構文でPL/CUDA関数の返り値に応じたデータを返します。

#plcuda_end

コードブロックの終了を宣言します。 なお、あるコードブロックの内側で他のコードブロックの開始を宣言した場合、現在のコードブロックは暗黙のうちに#plcuda_endディレクティブによって終了したものとして扱われます。

#plcuda_include <function name>

このディレクティブはCUDA Cの#includeと似ていますが、ヘッダファイルではなく、指定されたSQL関数の実行結果をディレクティブの存在していた場所に挿入します。 オプションで指定するSQL関数はPL/CUDA関数と同一の引数をとりtext型を返す必要があります。

これは例えば、大量データ間の類似度を計算する際に、計算のアルゴリズムはほとんど同一であるにも関わらず距離計算のロジックだけが異なるバリエーションを動的に作り出す事が可能で、PL/CUDA関数の保守を簡素化する事ができます。

#plcuda_library <library name>

CUDAプログラムをビルドする際にリンクするライブラリ名を指定します。 <library name>に記述するのは、nvccコマンドの-lオプションに相当する文字列です。 例えばlibcublas.soライブラリをリンクする場合には、接頭語のlibと拡張子の.soを省略したcublasと指定します。 現在のところ、CUDA Toolkitの標準ライブラリパス(/usr/local/cuda/lib64)にインストールされたライブラリのみを指定する事ができます。

#plcuda_sanity_check <function>

GPUカーネルの起動に先立って、引数の妥当性を検証するためのSQL関数をしています。 デフォルトでは妥当性検証関数は設定されていません。 GPUデバイスの初期化などを行う必要があるため、通常、GPUカーネル関数の起動はCPU上で別の関数を起動するよりも重い処理です。もし引数がPL/CUDA関数の仕様からは許容できない値を持っている場合、GPUカーネル関数を実行する数千~数百万(場合によってはそれ以上の)のGPUスレッドは、ただ引数の妥当性をチェックしてエラー状態を返却するためだけに起動されます。GPUカーネル関数を実行する前に、引数の妥当性チェックを十分小さなコストで行えるならば、妥当性検証関数を使用してGPUカーネル関数の実行前にエラーを発生させることを考慮すべきです。 妥当性検証関数は、PL/CUDA関数と同じ型の引数を持ち、bool返す関数です。

PL/CUDA 関連関数

関数定義 結果型 説明
plcuda_function_source(regproc) text 引数としてPL/CUDA関数のOIDを与えると、PL/CUDA関数から生成されるGPUカーネルのソースコードを返します。

PL/CUDA関数呼び出し支援

以下の関数群は、PL/CUDA関数の呼び出しを簡便にするために提供されています。

関数定義 結果型 説明
attnums_of(regclass,text[]) smallint[] 第一引数で指定したテーブルの第二引数で指定した列名(複数可)の列番号を配列として返します。
attnum_of(regclass,text) smallint 第一引数で指定したテーブルの第二引数で指定した列名の列番号を返します。
atttypes_of(regclass,text[]) regtype[] 第一引数で指定したテーブルの第二引数で指定した列名(複数可)のデータ型を配列として返します。
atttype_of(regclass,text) regtype 第一引数で指定したテーブルの第二引数で指定した列名のデータ型を返します。
attrs_types_check(regclass,text[],regtype[]) bool 第一引数で指定したテーブルの、第二引数で指定した列名(複数可)のデータ型が、第三引数で指定したデータ型とそれぞれ一致しているかどうかを調べます。
attrs_type_check(regclass,text[],regtype) bool 第一引数で指定したテーブルの、第二引数で指定した列名(複数可)のデータ型が、全て第三引数で指定したデータ型と一致しているかどうかを調べます。

配列ベースの行列型関数

本節ではPG-Stromの提供する配列ベースの行列型をサポートするSQL関数について説明します。

PostgreSQLには行列を表現するための専用のデータ型は存在していませんが、以下の条件を満たす二次元配列をあたかも行列であるかのように取り扱う事が可能です。

  • 二次元配列である
  • 各次元の配列要素が1から始まる
  • NULL値を含まない
  • 配列の大きさが1GBを越えない。(PostgreSQL可変長データ表現による制約)
  • smallintintbigintrealまたはfloat型の配列である

配列がこれらの条件を満たす時、行列の(i,j)要素の位置は添え字から一意に特定する事ができ、GPUスレッドが自らの処理すべきデータを効率的に取り出す事を可能とします。また、通常の行形式データとは異なり、計算に必要なデータのみをロードする事になるため、メモリ消費やデータ転送の点で有利です。 PG-Stromは、この様な疑似的な行列型をサポートするため、以下に示すSQL関数を提供しています。

関数定義 結果型 説明
array_matrix(variadic arg, ...) array 入力された行を全て連結した配列ベース行列を返す集約関数です。例えば、float型の引数x、y、zを1000行入力すると、同じfloat型で3列×1000行の配列ベース行列を返します。
この関数は可変長引数を取るよう定義されており、argは1個以上のsmallintintbigintrealまたはfloat型のスカラー値で、全てのarg値は同じデータ型を持つ必要があります。
matrix_unnest(array) record 配列ベース行列を行の集合に展開する集合関数です。arraysmallintintbigintrealまたはfloat型の配列で、行列の幅に応じて1個以上のカラムからなるrecord型を返却します。例えば、10列×500行から成る行列の場合、各レコードは行列要素のデータ型を持つ10個のカラムからなり、これが500行生成されます。
標準のunnest関数と似ていますが、record型を生成するため、AS (colname1 type[, ...])句を用いて返却されるべきレコードの型を指定する必要があります。
rbind(array, array) array arraysmallintintbigintrealまたはfloat型の配列です。
二つの配列ベース行列を縦方向に結合します。双方の行列は同一の要素データ型を持つ必要があり、行列の幅が等しくない場合は足りない部分を0で埋めます。
rbind(array) array arraysmallintintbigintrealまたはfloat型の配列です。rbind(array, array)と似ていますが、集合関数として動作し入力された全ての配列ベース行列を縦方向に結合します。
cbind(array, array) array arraysmallintintbigintrealまたはfloat型の配列で、二つの配列ベース行列を横方向に結合します。双方の行列は同一の要素データ型を持つ必要があり、行列の高さ等しくない場合は足りない部分を0で埋めます。
cbind(array) array arraysmallintintbigintrealまたはfloat型の配列で、cbind(array, array)と似ていますが、集合関数として動作し入力された全ての配列ベース行列を横方向に結合します。
transpose(array) array arraysmallintintbigintrealまたはfloat型の配列で、行列の幅と高さが入れ替わった転置行列を生成します。
array_matrix_validation(anyarray) bool 入力された配列(anyarray)が、配列ベース行列として妥当かどうかを検査します。 PL/CUDA関数実行前の引数の妥当性検証や、DOMAIN型を定義する時の検査制約としての利用を想定しています。
array_matrix_height(array) int arraysmallintintbigintrealまたはfloat型の配列で、配列ベース行列の高さを返却します。
array_matrix_width(array) int arraysmallintintbigintrealまたはfloat型の配列で、配列ベース行列の幅を返却します。