変数をとりあえず先に書きますが、APIの説明がないと
ほとんど意味が分からないので読み飛ばしましょう。
cl_platform_id platform_id = NULL;
cl_uint ret_num_platforms;
cl_device_id device_id = NULL;
cl_uint ret_num_devices;
cl_context context = NULL;
cl_command_queue command_queue = NULL;
cl_mem resmem = NULL;
cl_mem ope1mem = NULL;
cl_mem ope2mem = NULL;
cl_program program = NULL;
cl_kernel kernel = NULL;
size_t global_item_size = 4;
size_t local_item_size = 1;
cl_int ret;
6.カーネルファイルの読み込み(カーネルコードをホストコードに読み込ませる)
FILE *fp;
char fileName[] = "./add.cl";
char *source_str;
size_t source_size;
fp = fopen(fileName, "r");
if (!fp){
fprintf(stderr, "Failed to load kernel.\n");
exit(1);
}
source_str = (char*)malloc(MAX_SOURCE_SIZE);
source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp);
fclose(fp);
add.clを読み込みます。APIは使いません。
6番目の手順が最初にきて違和感がありますが、参考書の仕様です（
まあ、読み込むだけだしいつでも良いとは思います。C言語は関数の途中で
変数を宣言できないし、最初に持ってこないとややこしいとか？
1.プラットフォームの特定(動作環境)
ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);
使用可能なプラットフォームを取得します。
引数1：取得したいプラットフォームの個数で通常は1を指定します。
&platform_id：実際に取得できたプラットフォームのリストが格納される
&ret_num_platforms：実際に取得できたプラットフォームの個数が格納される
戻り値ret：APIの実行が成功したか失敗したか判断するエラーコードが格納される
※以後retに関してはこの用途で使われるので省略します。
2.デバイスの特定(nコアCPUを使うとかGPUを使うとか)
ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &ret_num_devices);
指定したプラットフォームで利用可能なデバイスを調べます。
引数platform_id：調べたいプラットフォームを指定する
CL_DEVICE_TYPE_DEFAULT：利用したいデバイスの種類を指定する
1：利用したいデバイスの個数を指定する
&device_id：実際に利用できるデバイスのリストが格納される
&ret_num_devices：実際に利用できるデバイスの個数が格納される
※CL_DEVICE_TYPE_DEFAULT：プラットフォームの標準デバイス
※CL_DEVICE_TYPE_CPU：ホスト側のCPU
※CL_DEVICE_TYPE_GPU：GPU
3.コンテキストの作成(GPUをいくつ使うとか実行環境)
context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);
コンテキストを作成します。これから作成する各種オブジェクトはこのコンテキストに所属することになります。
引数NULL：コンテキストのプロパティリスト？
1：利用したいデバイスの個数を指定する
&device_id：デバイスを指定する
NULL：コールバック関数？
NULL：コールバック関数の引数を指定する
戻り値context：作成されたコンテキストが格納される
4.コマンドキューの作成(コンテキストとかをぶっこんでホストとデバイスのやりとりを準備)
command_queue = clCreateCommandQueue(context, device_id, 0, &ret);
コマンドキューを作成します。このコマンドキューを通してホストからデバイスへの働きかけを決定します。
引数context：コマンドキューが所属するコンテキストを指定する
device_id：コマンドキューによる制御の対象となるデバイスを指定する
0：コマンドキュープロパティ？
戻り値command_queue：作成されたコマンドキューが格納される
5.メモリオブジェクトの作成(デバイス上で使用するメモリを確保)
ope1mem = clCreateBuffer(context, CL_MEM_READ_WRITE, ARRAY_SIZE*sizeof(float), NULL, &ret);
デバイス上にあるメモリを確保します。
引数context：メモリオブジェクトが所属するコンテキストを指定する
CL_MEM_READ_WRITE：メモリへのアクセス制限を指定する
ARRAY_SIZE*sizeof(float)：確保するメモリのサイズを指定する
NULL：ホスト側メモリポインタ？
戻り値ope1mem：メモリオブジェクトが格納される？
※他にも以下のようなアクセス制限がある
※CL_MEM_WRITE_ONLY：カーネル側からは書き込み専用
※CL_MEM_READ_ONLY：カーネル側からは読み取り専用
※CL_MEM_USR_HOST_PTR：ホスト側メモリポインタで指定された領域を使う
※CL_MEM_ALLOC_HOST_PTR：ホスト側からアクセス可能な領域を使う
※CL_MEM_COPY_HOST_PTR：ホスト側メモリポインタで指定された領域をコピー
※CL_MEM_HOST_WRITE_ONLY：ホスト側からは書き込み専用
※CL_MEM_HOST_READ_ONLY：ホスト側からは読み取り専用
※CL_MEM_HOST_NO_ACCRESS：ホスト側からは読み書き不可
メモリオブジェクトへのデータ転送（デバイス上で使用するメモリに値を設定する）
ret = clEnqueueWriteBuffer(command_queue, ope1mem, CL_TRUE, 0, ARRAY_SIZE*sizeof(float), ope1, 0, NULL, NULL);
デバイス上にあるメモリへデータをコピーします。
引数command_queue：コマンドキューを指定する
ope1mem：転送先のメモリオブジェクトを指定する
CL_TRUE：ブロッキング動作？の指定
0：コピー先メモリの先頭位置からのオフセットを指定する
ARRAY_SIZE*sizeof(float)：コピーするデータのサイズを指定する
ope1：コピー元のメモリポインタを指定する
0：実行前に終了を待つイベント？の数
NULL：実行前に終了を待つイベント？のリスト
NULL：この書き込みコマンドを識別するイベントオブジェクトが格納される
※CL_TRUE：書き込みが終了するまでこのAPIは終了しない（同期メモリコピー）
※CL_FALSE：コマンドキューへのキューイングが終了したらこのAPIを終了する（非同期メモリコピー）
7.プログラムオブジェクトの作成(読み込んだカーネルコードなどをプログラムオブジェクトとしておく)
program = clCreateProgramWithSource(context, 1, (const char **)&source_str, (const size_t *)&source_size, &ret);
プログラムオブジェクトを作成します。
引数context：プログラムオブジェクトは所属するコンテキストを指定する
1：カーネルコードを構成する文字列の個数
(const char **)&source_str：読み込んだカーネルコードを指定する
(const size_t *)&source_size：カーネルコードのサイズを指定
戻り値program：プログラムオブジェクトが格納される
8.ビルド(プログラムオブジェクトをコンパイル・リンクする)
ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
カーネルコードを実行可能にします。
引数program：対象となるプログラムオブジェクトを指定する
1：デバイスの個数を指定する
&device_id：デバイスのリストを指定する
NULL：ビルドオプション？
NULL：コールバック関数？
NULL：コールバック関数の引数を指定する
※コールバック関数がNULLの場合はビルド完了までこのAPIは終了しない
9.カーネルオブジェクトの作成(プログラムオブジェクトに含まれている1つ1つのカーネル関数が、1つ1つのカーネルオブジェクトになっているので抽出)
kernel = clCreateKernel(program, "add", &ret);
カーネルオブジェクトを作成します。
引数program：カーネルオブジェクトを含んだプログラムオブジェクトを指定する
"add"：カーネルオブジェクト化させるカーネル関数名を指定する
戻り値kernel：カーネルオブジェクトが格納される
10.カーネル引数の設定(カーネルオブジェクトは元カーネル関数なので引数を指定)
ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &ope1mem);
カーネル関数に引数を渡します。
引数kernel：カーネルオブジェクトを指定する
1：引数のインデックスを指定する（1なら第2引数に値を渡す）
sizeof(cl_mem)：渡す引数のサイズを指定する
(
void *) &ope1mem：渡す引数を指定する
11.カーネルの実行(コマンドキューにカーネルオブジェクトをぶっこむ)
ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global_item_size, &local_item_size, 0, NULL, NULL);
カーネルオブジェクトを実行します。
引数command_queue：デバイスを制御するコマンドキューを指定します。
kernel：実行するカーネルオブジェクトを指定します
1：ワークアイテム（処理単位）の次元？
NULL：ワークアイテムのグローバルIDを算出する際に使うオフセット値を指定する（NULLを指定すると各次元で0）
&global_item_size：各次元における全体のワークアイテムの総数を指定する
&local_item_size：ワークグループあたりのワークアイテムの個数を指定する
0：実行前に終了を待つイベントの数を指定する
NULL：実行前に終了を待つイベントのリストを指定する
NULL：この処理を表すイベント？が格納される
12.メモリオブジェクトからの読み込み(実行結果をデバイス上のメモリから入手)
ret = clEnqueueReadBuffer(command_queue, resmem, CL_TRUE, 0, ARRAY_SIZE*sizeof(float), result, 0, NULL, NULL);
デバイス上のメモリのデータをホスト上のメモリにコピーします。
引数command_queue：コマンドキューを指定します
resmem：コピー元のデバイス上のメモリを指定します
CL_TRUE：ブロッキング動作を指定します
0：コピー元メモリのオフセット
ARRAY_SIZE*sizeof(float)：コピーするデータのサイズを指定する
result：コピー先のホスト上のメモリを指定します
0：実行前に終了を待つイベントの数を指定する
NULL：実行前に終了を待つイベントのリストを指定する
NULL：この読み出しコマンドを識別するイベントオブジェクト？が格納される
13.オブジェクト解放(確保したメモリなどを解放)
ret = clFlush(command_queue);
ret = clFinish(command_queue);
ret = clReleaseKernel(kernel);
ret = clReleaseProgram(program);
ret = clReleaseMemObject(resmem);
ret = clReleaseMemObject(ope1mem);
ret = clReleaseMemObject(ope2mem);
ret = clReleaseCommandQueue(command_queue);
ret = clReleaseContext(context);
見たままなので説明省略します。
add.cl
__kernel void add(__global float *result, __global float *ope1, __global float *ope2){
int id = get_global_id(0);
result[id] = ope1[id] + ope2[id];
}
関数名に__kernelをつけ、仮引数に__globalをつければカーネル関数になるらしい。
あとはC言語に似ている書き方でOKみたいですが、それは数年後に必要になったら調べます。
ワークアイテムなどに関して
ワークアイテム：処理の最小単位。別の言語などだと「スレッド」のことだったり名前が違うかも？
ワークグループ：ワークアイテムを一定数まとめたもの。CUDAだと「ブロック」とか「グリッド」とかだけど若干違う気がする。
get_global_id(0)に関してですが、例えばワークアイテム12個を次のように管理した場合、
[■■■■][■■■■][
■■■■]
■：ワークアイテム
[ ]：ワークグループ
■は
2番目のワークグループ内の、
0番目のワークアイテムで、
全体として8番目のワークアイテム
であるのでカーネル側でこれらの数値を取得できます。その方法がget_global_id(0)
などの関数です。
get_group_id(0)：2
get_local_id(0)：0
get_global_id(0)：8
なお、引数の0はワークアイテムを多次元管理する場合に1なども使います。
PR
