PL/CUDA

This chapter introduces the way to implement GPU executable native program as SQL functions, using PL/CUDA procedural language. 本章では、PL/CUDA言語を用いて、GPUで実行可能なネイティブプログラムをSQL関数として実装する方法について説明します。

A series of discussion to support matrix data type is held in PostgreSQL developers' community now. Please note that the array-based matrix data type provided by the current version of PG-Strom may change its data format in the future release. 現在、PostgreSQLコミュニティでは行列型(matrix型)のサポートについて議論が行われています。 現バージョンのPG-Stromが独自に提供する行列型(配列を用いたmatrix型)は、将来のバージョンにおいてフォーマットが変更となる可能性がある事に留意してください。

PL/CUDA Overview PL/CUDA 概要

PG-Strom internally constructs GPU programs by CUDA language, according to the supplied SQL, then generates GPU's native binary using just-in-time compile. CUDA is a programming environment provided by NVIDIA. It allows implementing parallel program which is executable on GPU device, using C-like statement. The transformation process from SQL statement to CUDA program is internal process, thus, no need to pay attention what GPU programs are generated and executed from the standpoint of users. 内部的に、PG-StromはSQL構文を元にCUDA言語によるGPUプログラムを生成し、これを実行時コンパイルによってGPU用命令バイナリを生成します。 CUDAとはNVIDIA社の提供するプログラミング環境で、C言語に似た構文を用いてGPUで並列実行可能なプログラムを記述する事ができます。 SQL構文からCUDAプログラムへの変換プロセスは内部的なもので、ユーザの視点からは、どのようなGPU用プラグラムが生成、実行されるのかを意識する必要はありません。

On the other hands, PostgreSQL supports to add programming language to implement SQL functions by CREATE LANGUAGE statement. PL/CUDA is a language handler which supports CREATE LANGUAGE statement. It also allows to run arbitrary GPU programs implemented by users as SQL functions, but not only GPU programs automatically generated by PG-Strom based on SQL. 一方、PostgreSQLではCREATE LANGUAGE構文を用いてSQL関数の記述に用いるプログラミング言語を追加する事ができます。 PL/CUDAとはCREATE LANGUAGE構文に対応した言語ハンドラで、SQLを元にPG-Stromが自動生成するGPUプログラムだけでなく、ユーザが実装した任意のGPUプログラムをSQL関数として実行する事が可能となります。

Its argument can take the data types supported by PG-Strom, like numeric, text, or matrix data type. These arguments are implicitly transformed to GPU side by the PL/CUDA infrastructure, so users don't need to pay attention to data loading between the database and GPU devices. In a similar fashion, the return value of PL/CUDA function (including the case of variable length data type) will be written back to CPU from GPU, then decode to the result of SQL function. SQL関数の引数には、数値型やtext型、行列型など、PG-Stromのサポートするデータ型を使用する事ができますが、これらはPL/CUDA実行系が自動的にGPU側へデータを転送するため、データベースとGPU間のデータロードについて意識する必要はありません。また同様に、PL/CUDA関数の戻り値(可変長データ型である場合を含む)もGPU側からCPU側へと書き戻され、SQL関数の戻り値として整形されます。

Therefore, users can focus on productive tasks like implementation of statistical analysis, code optimization and so on, without routine process like data input/output between GPU and databases. そのため、ユーザはGPUやデータベースとの間のデータの入出力といった定型的な処理に煩わされる事なく、統計解析ロジックの実装や高速化といった生産的な作業に注力する事ができます。

Once a PL/CUDA function is declared with CREATE FUNCTION statement, it generates a CUDA program that embeds the definition of this function on the GPU's kernel function at the execution time. This kernel function contains initialization code to reference this PL/CUDA functions and auxiliary code to return run-time error to CPU side. Also, it can include some run-time functions to support execution of PG-Strom. CREATE FUNCTION構文を用いてPL/CUDA関数を定義すると、この関数の実行時、関数の定義部をそのままGPUのカーネル関数に埋め込んだCUDAプログラムを作成します。 このカーネル関数は、ユーザ定義処理の他に、PL/CUDA関数の引数を参照するための変数の初期化や、実行時エラーをCPU側へ返却するための補助的なコードを含んでいます。また、PG-Stromの実行をサポートするための各種ランタイム関数をインクルードする事もできます。

Here is no special memory protection mechanism on the native CUDA program made with PL/CUDA function, thus, execution of buggy PL/CUDA function can crash GPU execution environment or PostreSQL infrastructure in some cases. Thus, only database superuser can define PL/CUDA function. PL/CUDA関数を用いて作成したネイティブのCUDAプログラムには、特別なメモリ保護などの仕組みはなく、バグのあるPL/CUDA関数の実行により、GPU実行環境や場合によってはPostgreSQL側をクラッシュさせる事も可能です。したがって、PL/CUDA関数の定義はデータベース特権ユーザに限定されています。

Below is an example of simple PL/CUDA function. This function takes two int arguments, and then returns the sum of them with int data type. 以下に単純なPL/CUDA関数の例を示します。 この関数は、int型の引数を二つ取り、その和をint型で返却します。

postgres=# CREATE FUNCTION gpu_add(int, int)
RETURNS int
AS $$
#plcuda_include "cuda_mathlib.h"
#plcuda_begin
  if (get_global_id() == 0)
    *retval = pgfn_int4pl(kcxt, arg1, arg2);
#plcuda_end
$$ LANGUAGE plcuda;
CREATE FUNCTION

The code block enclosed by #plcuda_begin and #plcuda_end is main portion of PL/CUDA function. This kernel function can reference the int type argument as arg1 and arg2 which are pg_int4_t variables, and can return the result values written on the region pointed by retval variable which is a pointer of pg_int4_t * data type, as result of PL/CUDA function. pgfn_int4pl() is a runtime function of PG-Strom, declared at cuda_mathlib.h, which adds two pg_int4_t variables. #plcuda_begin#plcuda_endで囲まれた部分が、PL/CUDA関数の本体部分です。 int型の引数はそれぞれ、pg_int4_t型の変数arg1arg2として参照する事ができ、pg_int4_t *型のポインタretvalの示す領域にセットしたデータが、PL/CUDA関数の実行結果としてCPU側に返却されます。 pgfn_int4pl()cuda_mathlib.hで定義されたPG-Stromのランタイム関数の一つで、pg_int4_t同士の加算を実行します。

Below is an example of execution of this PL/CUDA function. Its two integer arguments (100 and 200) were sent to GPU device, then it wrote back the calculated result (300) from the GPU device. As like normal SQL functions, PL/CUDA function can be used as a part of SQL expression. このPL/CUDA関数を実行すると、以下のように引数である100, 200という整数値をGPU側に送出し、計算結果である300という値をGPUから書き戻しています。通常のSQL関数と同様に、PL/CUDA関数を他のSQL式の一部として使用する事もできます。

postgres=# SELECT gpu_add(100,200);
 gpu_add
---------
     300
(1 row)

The plcuda_function_source function allows showing the source of kernel function generated by the PL/CUDA function. The code block enclosed by the comment: code by pl/cuda function is the portion injected from the declaration of PL/CUDA function. PL/CUDA関数を定義した結果、どのようなカーネル関数が生成されるのかを確認するにはplcuda_function_source関数を使用します。 コメント文code by pl/cuda functionで囲まれたブロックがPL/CUDA関数の定義部から挿入された部分です。

postgres=# SELECT plcuda_function_source('gpu_add'::regproc);
                    plcuda_function_source
---------------------------------------------------------------
 #include "cuda_common.h"                                     +
 #include "cuda_mathlib.h"                                    +
 #include "cuda_plcuda.h"                                     +
                                                              +
 STATIC_INLINE(void)                                          +
 __plcuda_gpu_add_main(kern_plcuda *kplcuda,                  +
               void *workbuf,                                 +
               void *results,                                 +
               kern_context *kcxt)                            +
 {                                                            +
   pg_int4_t *retval __attribute__ ((unused));                +
   pg_int4_t arg1 __attribute__((unused));                    +
   pg_int4_t arg2 __attribute__((unused));                    +
   assert(sizeof(*retval) <= sizeof(kplcuda->__retval));      +
   retval = (pg_int4_t *)kplcuda->__retval;                   +
   arg1 = pg_int4_param(kcxt,0);                              +
   arg2 = pg_int4_param(kcxt,1);                              +
                                                              +
   /* ---- code by pl/cuda function ---- */                   +
   if (get_global_id() == 0)                                  +
     *retval = pgfn_int4pl(kcxt, arg1, arg2);                 +
   /* ---- code by pl/cuda function ---- */                   +
 }                                                            +
                                                              +
 KERNEL_FUNCTION(void)                                        +
 plcuda_gpu_add_main(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_main_kernel,kparams);     +
   __plcuda_gpu_add_main(kplcuda, workbuf, results, &kcxt);   +
   kern_writeback_error_status(&kplcuda->kerror_main, kcxt.e);+
 }                                                            +
                                                              +

(1 row)

Structure of PL/CUDA PL/CUDAの構造

Function declaration with PL/CUDA is consists of several code blocks split by directives that begin from #plcuda_. Only the code block start with #plcuda_begin is the minimum requirement, and you can add some other code block on demand. PL/CUDAの関数定義は、#plcuda_で始まるディレクティブによって分割されるいくつかのコードブロックから構成されます。 このうち、#plcuda_beginより始まるコードブロックのみが必須で、必要に応じてその他のコードブロックを追加する事ができます。

#plcuda_decl
  [...any declarations...]
#plcuda_prep
  [...function body of prep kernel...]
#plcuda_begin
  [...function body of main kernel...]
#plcuda_post
  [...function body of post kernel...]
#plcuda_end

The declaration block, which begins with #plcuda_decl, can have declaration of static functions we can call from other code blocks. Unlike other code blocks, the contents of the code block won't be injected into a particular kernel function, and you need to declare complete static functions. #plcuda_declより始まる宣言ブロックは、その他のコードブロックから呼び出す事ができるstatic関数の宣言を記述する事ができます。 他のコードブロックのように、コードブロックの内容が暗黙のうちに特定のカーネル関数に組み込まれる訳ではなく、完全な形式のstatic関数を定義する必要があります。 When a kernel function is executed with parallel threads larger than block size on a GPU device, the only way to synchronize between multiple execution units is synchronization of kernel function exit. For example, in case when algorithm is implemented under the assumption of correct initialization of the result buffer, you have to initialize the results buffer first, then you cannot execute the core of algorithm until completion of the initialization. If a part of threads would be executed towards uninitialized buffer, it easily leads incorrect calculation results or crash of execution environment, you always need to avoid. GPU上であるカーネル関数がブロックサイズを越える数のスレッドで並列実行されている時、複数の実行ユニット間で同期を取るには、カーネル関数終了のタイミングで待ち合わせる事が唯一の方法です。 例えば、結果バッファが特定の値で初期化されている事を前提としてアルゴリズムが実装されている場合、先ず、結果バッファの初期化を行い、それが全て完了するまではアルゴリズムの中核部分を実行する事はできません。 一部のスレッドが未初期化のバッファに対して実行されるという状況は、容易に不正確な計算結果や実行環境のクラッシュを招いてしまうため、常に避ける必要があります。

Every content of user defined code blocks, the preparation block begins from #plcuda_prep, the main block begins from #plcuda_begin, and the post-process block begins from #plcuda_post, shall be injected to the relevant kernel functions. Even though implementation of the preparation block and the post-process block are optional, we will ensure the order to launch the preparation kernel function, the main kernel function, then the post-process kernel function when these code blocks are defined. We intend to use these functions to initialize the results buffer or working buffer prior to execution of the main kernel function, or to summarize the final results next to execution of the main kernel. #plcuda_prepから始まる前処理ブロック、#plcuda_mainから始まる本体ブロック、および#plcuda_postから始まる後処理ブロックは、それぞれユーザ定義のコードブロックの内容が対応するカーネル関数に埋め込まれます。 前処理ブロックと後処理ブロックの定義はオプショナルですが、これらのコードブロックが定義されている時、前処理カーネル関数、本体カーネル関数、後処理カーネル関数の順で実行される事が保証されています。 これらは、本体カーネル関数の実行に先立って結果バッファや作業バッファの初期化を行う事や、本体カーネル関数の実行後に最終結果を集計するなどの用途に使用する事を意図しています。

An invocation of PL/CUDA function internall contains several SQL functions and launch GPU kernel functions. Prior to the GPU kernel functions, we have to determine the parameters when GPU kernel functions like number of threads, amount of results and working buffer. These parameters depend on the arguments, so PL/CUDA handler determines with other SQL functions that take identical argument signature. 一個のPL/CUDA関数の呼び出しは、内部的には何個かのSQL関数、GPUカーネル関数の呼び出しを含んでいます。 GPUカーネル関数の呼び出しに先立って、GPUカーネル関数を起動する際のスレッド数、作業バッファや結果バッファのサイズといったパラメータを決定する必要があります。 これらは引数により変動するため、PL/CUDA言語ハンドラは、同じ引数を取る他のSQL関数を呼びだしてこれらのパラメータを決定します。

Once we could determine the parameters to call GPU kernel function, PL/CUDA handler transfers the arguments of PL/CUDA function to the argument buffer on GPUs, by DMA copy. GPUカーネル関数の呼び出しパラメータが確定すると、次に、PL/CUDA言語ハンドラは、DMAを用いてPL/CUDA関数の引数をGPU上の引数バッファに転送します。

Then, it launches the preparation kernel function (if any), the main kernel function, and the post-process kernel function (if any). Please note that we cannot synchronize GPU threads across the block size boundary, except for the timing of GPU kernel function begin/end. It means, if you expect a particular state exists on the working buffer or results buffer, buffer initialization by preparation kernel then reference of this data structure by the main kernel are required. 続いて、(定義されていれば)前処理カーネル関数、本体カーネル関数、(定義されていれば)後処理カーネル関数を呼びだします。ブロックサイズを越えたGPUスレッド間で同期を取る方法は、GPUカーネル関数の開始終了のタイミング以外に無い事に留意してください。つまり、作業バッファや結果バッファがある特定の状態を持っている事を期待するのであれば、前処理カーネル関数で初期化を行い、次に本体カーネル関数でこれらのデータ構造を参照する必要があります。

Finally, PL/CUDA handler writes back the contents of result buffer into the host side. In case when PL/CUDA function returns a fixed-length datum, the code block updates the area pointed by the retval variable which is initialized prior to execution of the user's defined block. In case when PL/CUDA function returns a variable-length datum, retval points to the area of pg_varlena_t, and its value has to be a reference to the results buffer (void *results), if it is not a NULL. Please note that it shall not be written back if retval points out of the results buffer. 最後に、PL/CUDA言語ハンドラは結果バッファの内容を本体側へ書き戻します。 PL/CUDA関数が固定長のデータを返す場合、GPUカーネル関数がユーザ定義ブロックの開始前に設定する変数retvalポインタの示す領域を更新します。 PL/CUDA関数が可変長のデータを返す場合、retvalpg_varlena_t型の領域を指しており、その値が非NULLである場合には結果バッファ(void *results)への参照でなければいけません。結果バッファ以外の領域を指していたとしても、これは本体側へ書き戻されない事に留意してください。

typedef struct {
    varlena    *value;      /* reference to the results buffer */
    cl_bool     isnull;     /* true, if NULL */
} pg_varlena_t;

#plcuda_num_threads directive allows specifying the number of threads to execute GPU kernel function. This directive can be used inside of the code block, and takes either a constant value or a SQL function. This SQL function has to be declared to take identical argument types and return bigint type. GPUカーネル関数を実行するスレッド数を指定するには#plcuda_num_threadsディレクティブを使用します。このディレクティブはコードブロックの内側で使用され、定数値またはSQL関数名を指定します。SQL関数は、PL/CUDA関数と同一の引数を持ちbigint型を返す関数として宣言されている必要があります。

In a similar fashion, #plcuda_shmem_unitsz allows to specify the amount of shared memory per thread, to be acquired on GPU kernel function launch. For example, when a GPU kernel function that consumes 8bytes per thread is launched with 384 threads per streaming-multiprocessor, 3KB of shared memory shall be available. Please note that the number of threads per streaming-multiprocessor shall be automatically calculated during the code optimization, a different concept from what we specify with #plcuda_num_threads directive. 同様に、#plcuda_shmem_unitszディレクティブを使用する事で、GPUカーネル関数の実行時に動的に確保する共有メモリのサイズを、スレッドあたりの大きさで指定する事ができます。例えば、スレッドあたり8バイトの共有メモリを使用するGPUカーネル関数が実行ユニットあたり384スレッドで起動された場合、3KBの共有メモリを使用する事ができます。 ここで言う実行ユニットあたりスレッド数は、最適化の結果自動的に算出される値で、#plcuda_num_threadsで指定する値とは異なる事に留意してください。

#plcuda_kernel_maxthreads directive allows switching optimization policy of the kernel function for the current code block, from maximization of execution efficiency to maximization of number of threads per streaming-multiprocessor (usually 1024). Increase of number of threads per streaming-multiprocessor will improve the performance of workloads which heavily use inter-threads synchronization using shared memory, like reduction operation. また、#plcuda_kernel_maxthreadsディレクティブを使用する事で、コードブロックから作成されるカーネル関数の最適化方針を、実行効率最大化から、実行ユニットあたりスレッド数最大化(通常、1024スレッド)へと切り替える事が可能です。実行ユニットあたりのスレッド数が増加する事で、縮約演算など、共有メモリを用いた実行ユニット間の同期処理を中核とする処理での性能向上か期待できます。

#plcuda_num_threads (<value>|<function name>)
#plcuda_shmem_unitsz  (<value>|<function name>)
#plcuda_kernel_maxthreads

PL/CUDA Reference PL/CUDAリファレンス

PL/CUDA Directives PL/CUDA ディレクティブ

This section introduces the directives we can use in PL/CUDA functions. 本節ではPL/CUDA関数で使用する事のできるディレクティブについて説明します。

#plcuda_begin

It marks beginning of the main kernel function code block. This directive is always required. 本体カーネル関数のコードブロックの開始を宣言します。このディレクティブは必須です。 Prior to execution of the code block on GPU, the arguments of PL/CUDA function are initialized for references by variable names like arg1, arg2, ... These variables have same representation with what PG-Strom represents SQL data types on GPU, for example, an argument of the real data type (that is single precision floating point type) is shown as a pg_float4_t type variable as declared below. GPU上でのコードブロックの実行開始に先立って、PL/CUDA関数の引数はarg1arg2、...という変数名で参照可能となるよう初期化されます。 これらの変数は、PG-StromがSQLデータ型をGPU上で表現するのと同じ表現を持っており、例えば、単精度浮動小数点であるreal型の引数は、以下のように定義されたpg_float4_t型の変数として表現されています。

typedef struct {
    cl_float    value;
    cl_bool     isnull;
} pg_float4_t;

These variables are kept in private area of each threads, thus, update of these variables are not reflected on execution of the kernel function on the next step. If you want to share the state between kernel functions, value shall be kept in either the working buffer referenced by the void *workbuf pointer or the results buffer referenced by the void *results pointer. これらの変数は各スレッドのプライベート領域に確保されており、変数を更新したとしても次ステップのカーネル関数には反映されません。カーネル関数の終了後、次のカーネル関数に状態を引き継ぐには、void *workbufポインタが参照する作業バッファか、void *resultsポインタの参照する結果バッファに値を格納する必要があります。

#plcuda_end
It marks end of the kernel function code block. By the way, if a directive to start code block was put inside of the different code block, the current code block is implicitly closed by the #plcuda_end directive. コードブロックの終了を宣言します。 なお、あるコードブロックの内側で他のコードブロックの開始を宣言した場合、現在のコードブロックは暗黙のうちに#plcuda_endディレクティブによって終了したものとして扱われます。
#plcuda_decl
Use of this directive is optional. このディレクティブの使用は任意です。 It marks beginning of the declaration code block that contains the raw code to be declared prior to the definition of any kernel functions. Unlike other code blocks, the contents of this code block shall not be applied as a kernel function, thus, you have to put complete definition of functions. 全てのkernel関数の定義に先立って宣言しておくべきコードブロックの開始を宣言します。 他のコードブロックとは異なり、内容が自動的にkernel関数として展開される事はありませんので、完全な関数定義を記述する必要があります。
#plcuda_prep
Use of this directive is optional. このディレクティブの使用は任意です。 It marks beginning of the preparation code block that shall be executed on GPU prior to the main kernel function; begins from #plcuda_begin directive. We expect the preparation kernel initializes the results and working buffer. The main kernel shall not be kicked until completion of the preparation kernel. #plcuda_beginから始まる本体カーネル関数の実行に先立ってGPUで実行すべき、前処理カーネル関数の処理を記述します。 ここでは、結果バッファや作業バッファの初期化を行う事を意図しており、前処理カーネル関数の実行が完了するまでは本体カーネル関数は実行されません。 Arguments of PL/CUDA functions can be referenced like as the main kernel function doing. PL/CUDA関数の引数へは、本体カーネル関数と同様にアクセスする事ができます。
#plcuda_post
You can optionally use this directive. このディレクティブの使用は任意です。 It marks beginning of the post-process code block that shall be executed on GPU next to the main kernel function; begins from #plcuda_begin directive. We expect the post-process kernel set up the final results to be returned to the CPU side. The post-process kernel shall not be kicked until completion of the preparation kernel. #plcuda_beginから始まる本体カーネル関数の実行後にGPUで実行すべき、後処理カーネル関数の処理を記述します。 ここでは、CPU側に返却する最終結果を結果バッファにセットする事を意図しており、本体カーネル関数の実行が完了するまでは後処理カーネル関数は実行されません。 Arguments of PL/CUDA functions can be referenced like as the main kernel function doing. PL/CUDA関数の引数へは、本体カーネル関数と同様にアクセスする事ができます。
#plcuda_num_threads (<value>|<function>)
Use of this directive is optional. If not specified, the default is a constant value 1.

このディレクティブの使用は任意です。未指定の場合、デフォルト値として定数1が使われます。

This directive allows specifying the number of threads to execute the GPU kernel function if it is used in the code block of #plcuda_prep, #plcuda_begin, or #plcuda_post. このディレクティブが#plcuda_prep#plcuda_begin、および#plcuda_postコードブロックの内側で指定されると、それぞれのGPUカーネル関数を起動する際のスレッド数を指定する事ができます。

If a constant value is specified, PL/CUDA runtime kicks the specified number of GPU threads to run the GPU kernel function. If a SQL function name is specified, PL/CUDA runtime call the specified SQL function, and then result of the function shall be applied as the number of GPU threads to run the GPU kernel function. This SQL function takes identical arguments with PL/CUDA function, and returns bigint data type. 数値が指定されると、PL/CUDAランタイムは指定された数のGPUスレッドを起動してGPUカーネル関数を実行します。 関数名が指定されると、PL/CUDAランタイムは指定されたSQL関数を呼び出し、戻り値で指定された数のGPUスレッドを起動します。このSQL関数は、PL/CUDA関数と同一の引数を取り、bigint型を返す必要があります。

#plcuda_shmem_unitsz (<value>|<function>)

Use of this directive is optional. If not specified, the default is a constant value 0. このディレクティブの使用は任意です。未指定の場合のデフォルト値は定数0です

This directive allows specifying amount of the shared memory per thread to be dinamically allocated on GPU kernel execution, if it is used in the code block of #plcuda_prep, #plcuda_begin, or #plcuda_post. このディレクティブが#plcuda_prep#plcuda_begin、および#plcuda_postコードブロックの内側で指定されると、それぞれのGPUカーネル関数を起動する際に動的に確保するスレッドあたり共有メモリのサイズを指定する事ができます。

If a constant value is specified, PL/CUDA runtime kicks GPU kernel function with the specified amount of the shared memory per thread. If a SQL function name is specified, PL/CUDA runtime call the specified SQL function, and then result of the function shall be applied as the amount of the shared memory per thread to run the GPU kernel function. This SQL function takes identical arguments with PL/CUDA function, and returns bigint data type. 数値が指定されると、PL/CUDAランタイムは指定された大きさのスレッドあたり共有メモリを確保してGPUカーネル関数を実行します。 関数名が指定されると、PL/CUDAランタイムは指定されたSQL関数を呼び出し、戻り値で指定された大きさのスレッドあたり共有メモリを確保してGPUカーネル関数を実行します。このSQL関数は、PL/CUDA関数と同一の引数を取り、bigint型を返す必要があります。

Please note that amount of the shared memory actually acquired on execution of GPU kernel function depends on the number of threads per streaming-multiprocessor, not only the amount of shared memory per thread specified by this directive. (Also note that the number of threads per streaming-multiprocessor is a different concept what we specified using #plcuda_num_threads.) For example, if amount of shared memory per thread is 8 bytes and the number of streaming-multiprocessor is 384, 3KB of shared memory shall be allocated per streaming-multiprocessor. At that time, if the number of total threads specified by #plcuda_num_threads is 32768, this GPU kernel shall be executed with 86 streaming-multiprocessor. However, it is the role of scheduler to determine the timing to put kernels into, so it does not mean that 86 x 3KB = 256KB of the shared memory is consumed at once. GPUカーネル関数の実行時に実際に確保される共有メモリのサイズは、本ディレクティブによって指定したスレッドあたり共有メモリのサイズだけでなく、実行ユニットあたりのスレッド数に依存する事に留意してください。(また、実行ユニットあたりのスレッド数は#plcuda_num_threadsで指定した値とも異なる概念である事に留意してください。) 例えば、スレッドあたり共有メモリのサイズが8バイトであり、実行ユニットあたりのスレッド数が384である場合、実行ユニット毎に3KBの共有メモリが確保されます。この時、#plcuda_num_threadsで指定したスレッド数が32768であれば、このGPUカーネルは86個の実行ユニットを使用して実行されますが、実行ユニットにタスクが投入されるタイミングを決めるのはスケジューラの役割ですので、必ずしも3KB x 86個 = 258KBの共有メモリが一度に消費されるわけではありません。

#plcuda_shmem_blocksz (<value>|<function>)

Use of this directive is optional. If not specified, the default is a constant value 0. このディレクティブの使用は任意です。未指定の場合のデフォルト値は定数0です

This directive allows specifying amount of the shared memory per block to be dinamically allocated on GPU kernel execution, if it is used in the code block of #plcuda_prep, #plcuda_begin, or #plcuda_post. このディレクティブが#plcuda_prep#plcuda_begin、および#plcuda_postコードブロックの内側で指定されると、それぞれのGPUカーネル関数を起動する際に動的に確保するブロックあたり共有メモリのサイズを指定する事ができます。

If a constant value is specified, PL/CUDA runtime kicks GPU kernel function with the specified amount of the shared memory per block. If a SQL function name is specified, PL/CUDA runtime call the specified SQL function, and then result of the function shall be applied as the amount of the shared memory per block to run the GPU kernel function. This SQL function takes identical arguments with PL/CUDA function, and returns bigint data type. 数値が指定されると、PL/CUDAランタイムは指定された大きさのブロック毎共有メモリを確保してGPUカーネル関数を実行します。 関数名が指定されると、PL/CUDAランタイムは指定されたSQL関数を呼び出し、戻り値で指定された大きさのブロックあたり共有メモリを確保してGPUカーネル関数を実行します。このSQL関数は、PL/CUDA関数と同一の引数を取り、bigint型を返す必要があります。

#plcuda_kernel_blocksz (<value>|<function>)

Use of this directive is optional. このディレクティブの使用は任意です。

This directive allows specifying the number of threads per streaming-multiprocessor, if it is used in the code block of #plcuda_prep, #plcuda_begin, or #plcuda_post. It is usually a multiple number of the warp value of the device, and equal to or less than 1024. In the default, an optimal value is applied according to the resource consumption of the GPU kernel function, therefore, this directive shall not be used unless you have no special reason; a larger block size is preferable due to characteristics of the algorithm for example. #plcuda_prep#plcuda_begin、および#plcuda_postコードブロックの内側でこのディレクティブを指定すると、実行ユニットあたりのスレッド数を指定する事ができます。通常、この値はデバイスのWARP値の倍数であり1024以下です。 デフォルトでは、GPUカーネル関数のリソース消費量に基づいた最適な値が使用されます。したがって、アルゴリズムの性質上大きなブロックサイズが望ましいなどの理由がない限り、本ディレクティブを使用すべきではありません。

If a constant value is specified, PL/CUDA runtime kicks GPU kernel function with the specified amount of the shared memory per block. If a SQL function name is specified, PL/CUDA runtime call the specified SQL function, and then result of the function shall be applied as the amount of the shared memory per block to run the GPU kernel function. This SQL function takes identical arguments with PL/CUDA function, and returns bigint data type. 数値が指定されると、PL/CUDAランタイムは指定された大きさのブロックを設定してGPUカーネル関数を実行します。 関数名が指定されると、PL/CUDAランタイムは指定されたSQL関数を呼び出し、戻り値で指定された大きさのブロックを設定してGPUカーネル関数を実行します。このSQL関数は、PL/CUDA関数と同一の引数を取り、bigint型を返す必要があります。

Increase the number of threads per streaming-multiprocessor allows more threads to synchronize other threads using the shared memory, on the other hands, it leads decrease of the amount of registers a thread can use, thus, it may have performance degradation by private variables allocation on the (slow) global memory for example. ブロックあたりスレッド数が多くなると、より多くのスレッドが共有メモリを介して同期処理を行う事ができるようになる半面、スレッドが使用できるレジスタ数が減少するため、一部のローカル変数がグローバルメモリ上に確保されるなど性能面では不利になる事があります。

#plcuda_include "library name"
It allows including the GPU runtime functions of PG-Strom, and using by the PL/CUDA functions. Please note that it is not a feature to include arbitrary header file. PG-StromのGPUランタイム関数をインクルードし、PL/CUDA関数内で使用できるようにします。 任意のヘッダファイルをインクルードして利用するための機能ではない事に留意してください。
  • "cuda_dynpara.h"

    It is a collection of GPU runtime functions related to dynamic parallelism; that launch kernel functions on GPU. Include of this file also links the device runtime library of CUDA. GPU内で動的にカーネル関数を起動するDynamic Parallelism関連のGPUランタイム関数群です。 このファイルをインクルードすると、CUDAのデバイスランタイムも同時にリンクされるようになります。

  • "cuda_matrix.h"

    It is a collection of GPU runtime functions to process the matrix data type of SQL. Please note that the matrix data type in this version is provided by PG-Strom, thus, it may not be compatible when the future version of PostgreSQL newly supports the matrix data type. SQLの行列型を処理するためのGPUランタイム関数群です。 なお、現バージョンの行列型はPG-Stromが独自に提供しているものであり、将来バージョンのPostgreSQLが行列型に対応した際に、フォーマットの互換性が維持されない可能性に留意してください。

  • "cuda_timelib.h"

    It is a collection of GPU runtime functions to process the date and time data type of SQL. SQLの日付時刻型を処理するためのGPUランタイム関数群です。

  • "cuda_textlib.h"

    It is a collection of GPU runtime functions to process the text data type and LIKE operator. SQLのテキストデータ型、およびLIKEオペレータを処理するためのGPUランタイム関数群です。

  • "cuda_numeric.h"

    It is a collection of GPU runtime functions to process the numeric data type of SQL. SQLのNumericデータ型を処理するためのGPUランタイム関数群です。

  • "cuda_mathlib.h"

    It is a collection of GPU runtime functions to process the arithmetic operators and mathematic functions of SQL. SQLの数学関数や四則演算オペレータを処理するためのGPUランタイム関数群です。

  • "cuda_money.h"

    It is a collection of GPU runtime functions to process the currency data type of SQL. SQLの通貨型を処理するためのGPUランタイム関数群です。

  • #plcuda_results_bufsz (<value>|<function>)

    Use of this directive is optional. If not specified, the default is a constant value 0. このディレクティブの使用は任意です。未指定の場合のデフォルト値は定数0です

    This directive allows specifying amount of the results buffer in bytes, to be acquired on execution of PL/CUDA function. If PL/CUDA function is declared to return variable length datum, allocation of the results buffer is needed. PL/CUDA関数の実行時に確保する結果バッファの大きさをバイト単位で指定します。PL/CUDA関数が可変長型データを返却する際には、結果バッファの確保は必須です。

    If a constant value is specified, PL/CUDA language handler acquires the specified amount of GPU RAM as the results buffer, then launch the GPU kernel functions. If a SQL function name is specified, PL/CUDA language handler call the specified SQL function, then result of the function shall be applied as the amount of GPU RAM for the results buffer and launch the GPU kernel functions. This SQL function takes identical arguments with PL/CUDA function, and returns bigint data type. 数値が指定されると、PL/CUDA言語ハンドラは指定されたバイト数のGPU RAMを結果バッファとして確保してからGPUカーネル関数を起動します。 関数名が指定されると、PL/CUDA言語ハンドラは指定されたSQL関数を呼び出し、戻り値で指定されたバイト数のGPU RAMを結果バッファとして確保し、GPUカーネル関数を起動します。このSQL関数は、PL/CUDA関数と同一の引数を取り、bigint型を返す必要があります。

    GPU kernel functions can access the results buffer as the region pointed by the void *results argument. If 0 bytes were specified, NULL shall be set on the void *results. GPUカーネル関数からは、結果バッファは引数void *resultsで指定された領域としてアクセス可能です。 0バイトが指定された場合、void *resultsにはNULLがセットされます。

    #plcuda_working_bufsz (<value>|<function>)

    Use of this directive is optional. If not specified, the default is a constant value 0. このディレクティブの使用は任意です。未指定の場合のデフォルト値は定数0です

    This directive allows specifying amount of the working buffer in bytes, to be acquired on execution of PL/CUDA function. PL/CUDA関数の実行時に確保する作業バッファの大きさをバイト単位で指定します。

    If a constant value is specified, PL/CUDA language handler acquires the specified amount of GPU RAM as the working buffer, and then launch the GPU kernel functions. If a SQL function name is specified, PL/CUDA language handler call the specified SQL function, then result of the function shall be applied as the amount of GPU RAM for the working buffer and launch the GPU kernel functions. This SQL function takes identical arguments with PL/CUDA function, and returns bigint data type. bigint型を返す必要があります。

    GPU kernel functions can access the working buffer as the region pointed by the void *results argument. If 0 bytes were specified, NULL shall be set on the void *results. GPUカーネル関数からは、作業バッファは引数void *workbufで指定された領域としてアクセス可能です。 0バイトが指定された場合、void *workbufにはNULLがセットされます。

    #plcuda_sanity_checl <function>

    It allows to specify the sanity check function that preliminary checks adequacy of the supplied arguments, prior to GPU kernel launch. No sanity check function is configured on the default. GPUカーネルの起動に先立って、引数の妥当性を検証するためのSQL関数をしています。 デフォルトでは妥当性検証関数は設定されていません。 Usually, launch of GPU kernel function is heavier task than call of another function on CPU. If supplied arguments have unacceptable values from the specification of the PL/CUDA function, a few thousands or millions (or more in some cases) of GPU kernel threads shall be launched just to check the arguments and return an error status. If sanity check can be applied prior to the launch of GPU kernel function with enough small cost, it is a valuable idea to raise an error using sanity check function prior to the GPU kernel function. 通常、GPUカーネル関数の起動はCPU上で別の関数を起動するよりも重い処理です。 もし引数がPL/CUDA関数の仕様からは許容できない値を持っている場合、GPUカーネル関数を実行する数千~数百万(場合によってはそれ以上の)のGPUスレッドは、ただ引数の妥当性をチェックしてエラー状態を返却するためだけに起動されます。GPUカーネル関数を実行する前に、引数の妥当性チェックを十分小さなコストで行えるならば、妥当性検証関数を使用してGPUカーネル関数の実行前にエラーを発生させることを考慮すべきです。 The sanity check function takes identical arguments with PL/CUDA function, and returns bool data type. 妥当性検証関数は、PL/CUDA関数と同じ型の引数を持ち、bool返す関数です。

    #plcuda_cpu_fallback <function>

    It allows to specify the CPU fallback function that performs as like GPU kernel function. No CPU fallback function is configured on the default. GPUカーネル関数と同等の処理を行うCPUフォールバック関数を指定します。 デフォルトではCPUフォールバック関数は設定されていません。

    If GPU kernel function returns StromError_CpuReCheck error and the CPU fallback function is configured, the PL/CUDA language handler discards the results of processing on GPU side, then call the CPU fallback function. It is valuable to implement an alternative remedy, in case when GPU kernel function is not always executable for all possible input; for example, data size may be too large to load onto GPU RAM. Also note that we must have a trade-off of the performance because CPU fallback function shall be executed in CPU single thread. GPUカーネル関数がStromError_CpuReCheckエラーを返却し、さらにCPUフォールバック関数が設定されていると、PL/CUDA言語ハンドラはGPUでの処理結果を破棄してCPUフォールバック関数を呼びだします。 これは、必ずしも全ての入力に対してGPUカーネル関数を実行可能でない(例えばデータサイズがGPU RAMに載りきらないなど)場合に、代替の救済策を実装するために有用です。ただし、CPUフォールバック関数はシングルスレッドで実行されるため、パフォーマンスが犠牲にならざるを得ない点には留意してください。

    PL/CUDA related functions PL/CUDA関連関数

    This section introduces the SQL functions related to PL/CUDA. 本節ではPL/CUDAに関連するSQL関数について説明します。

    SQL Functions related to PL/CUDA PL/CUDA関連SQL関数
    Definition 関数定義 Return Value 戻り値 Description 説明
    plcuda_function_source(regproc) text It returns source code of the GPU kernel generated from the PL/CUDA function, towards the OID input of PL/CUDA function as argument. 引数としてPL/CUDA関数のOIDを与えると、PL/CUDA関数から生成されるGPUカーネルのソースコードを返します。

    Array-based matrix type 配列ベースの行列型

    This section introduces the SQL functions that supports array-based matrix types provided by PG-Strom. 本節ではPG-Stromの提供する配列ベースの行列型をサポートするSQL関数について説明します。

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

  • 2-dimensional Array 二次元配列である
  • Element of array begins from 1 for each dimension 各次元の配列要素が1から始まる
  • No NULL value is contained NULL値を含まない
  • Length of the array is less than 1GB, due to the restriction of variable length datum in PostgreSQL 配列の大きさが1GBを越えない。(PostgreSQL可変長データ表現による制約)
  • Array with smallint, int, bigint, real or float data type smallintintbigintrealまたはfloat型の配列である
  • If and when the array satisfies the above terms, we can determine the location of (i,j) element of the array by the index uniquely, and it enables GPU thread to fetch the datum to be processed very efficiently. Also, array-based matrix packs only the data to be used for calculation, unlike usual row-based format, so it has advantaged on memory consumption and data transfer. 配列がこれらの条件を満たす時、行列の(i,j)要素の位置は添え字から一意に特定する事ができ、GPUスレッドが自らの処理すべきデータを効率的に取り出す事を可能とします。また、通常の行形式データとは異なり、計算に必要なデータのみをロードする事になりますので、メモリ消費やデータ転送の点で有利です。 PG-Stromは、この様な疑似的な行列型をサポートするため、以下に示すSQL関数を提供しています。

    SQL Functions related to Array-based Matrix 配列ベース行列関連 SQL関数
    Definition 関数定義 Return Value 戻り値 Description 説明
    array_matrix(variadic arg, ...) array It is an aggregate function that combines all the rows supplied. For example, when 3 float arguments were supplied by 1000 rows, it returns an array-based matrix with 3 columns X 1000 rows, float data type. 入力された行を全て連結した配列ベース行列を返す集約関数です。例えば、float型の引数xyzを1000行入力すると、同じfloat型で3列×1000行の配列ベース行列を返します。
    This function is declared to take variable length arguments. The arg takes one or more scalar values of either smallint, int, bigint, real or float. All the arg must have same data types. この関数は可変長引数を取るよう定義されており、argは1個以上のsmallintintbigintrealまたはfloat型のスカラー値で、全てのarg値は同じデータ型を持つ必要があります。
    matrix_unnest(array) record It is a set function that extracts the array-based matrix to set of records. array is an array of smallint, int, bigint, real or float data. It returns record type which consists of more than one columns according to the width of matrix. For example, in case of a matrix with 10 columns X 500 rows, each records contains 10 of columns with element type of the matrix, then it generates 500 of the records. 配列ベース行列を行の集合に展開する集合関数です。arraysmallintintbigintrealまたはfloat型の配列で、行列の幅に応じて1個以上のカラムからなるrecord型を返却します。例えば、10列×500行から成る行列の場合、各レコードは行列要素のデータ型を持つ10個のカラムからなり、これが500行生成されます。
    It is similar to the standard unnest function, but generates record type, thus, it requires to specify the record type to be returned using AS (colname1 type[, ...]) clause. 標準のunnest関数と似ていますが、record型を生成するため、AS (colname1 type[, ...])句を用いて返却されるべきレコードの型を指定する必要があります。
    rbind(array, array) array array is an array of smallint, int, bigint, real or float data. arraysmallintintbigintrealまたはfloat型の配列です。
    This function combines the supplied two matrices vertically. Both matrices needs to have same element data type. If width of matrices are not equivalent, it fills up the padding area by zero. 二つの配列ベース行列を縦方向に結合します。双方の行列は同一の要素データ型を持つ必要があり、行列の幅が等しくない場合は足りない部分を0で埋めます。
    rbind(array) array array is an array of smallint, int, bigint, real or float data. This function is similar to rbind(array, array), but performs as an aggregate function, then combines all the input matrices into one result vertically. arraysmallintintbigintrealまたはfloat型の配列です。 rbind(array, array)と似ていますが、集合関数として動作し入力された全ての配列ベース行列を縦方向に結合します。
    cbind(array, array) array array is an array of smallint, int, bigint, real or float data. This function combines the supplied two matrices horizontally. Both matrices needs to have same element data type. If height of matrices are not equivalent, it fills up the padding area by zero. arraysmallintintbigintrealまたはfloat型の配列で、二つの配列ベース行列を横方向に結合します。双方の行列は同一の要素データ型を持つ必要があり、行列の高さ等しくない場合は足りない部分を0で埋めます。
    cbind(array) array array is an array of smallint, int, bigint, real or float data. This function is similar to rbind(array, array), but performs as an aggregate function, then combines all the input matrices into one result horizontally. arraysmallintintbigintrealまたはfloat型の配列で、rbind(array, array)と似ていますが、集合関数として動作し入力された全ての配列ベース行列を横方向に結合します。
    transpose(array) array array is an array of smallint, int, bigint, real or float data. This function makes a transposed matrix that swaps height and width of the supplied matrix. arraysmallintintbigintrealまたはfloat型の配列で、行列の幅と高さが入れ替わった転置行列を生成します。
    array_matrix_validation(anyarray) bool It validates whether the supplied array (anyarray) is adequate for the array-based matrix. It is intended to use for sanity check prior to invocation of PL/CUDA function, or check constraint on domain type definition. 入力された配列(anyarray)が、配列ベース行列として妥当かどうかを検査します。 PL/CUDA関数実行前の引数の妥当性検証や、DOMAIN型を定義する時の検査制約としての利用を想定しています。
    array_matrix_height(array) int array is an array of either smallint, int, bigint, real or float data. This function returns the height of the supplied matrix. arraysmallintintbigintrealまたはfloat型の配列で、行列の高さを返却します。
    array_matrix_width(array) int array is an array of either smallint, int, bigint, real or float data. This function returns the width of the supplied matrix. arraysmallintintbigintrealまたはfloat型の配列で、行列の幅を返却します。
    array_matrix_rawsize(regproc,int,int) bigint It returns the required length of memory to store an array-based matrix that has the data type specified by the 1st argument, height by the 2nd argument and width by the 3rd argument. It is intended to use for #plcuda_results_bufsz and #plcuda_working_bufsz. 第一引数で指定したデータ型、第二引数で指定した高さ、および第三引数で指定した幅の配列ベース行列の作成に必要なメモリ領域の大きさを返します。#plcuda_results_bufsz#plcuda_working_bufszでの利用を意図しています。