目指すところ
OpenCLの正しい使い方みたいな話は他のサイトに譲って、とりあえず使いましょう。楽しいから。
C++でOpenCL(環境構築編)もあります。
プログラム全文
githubからクローンできます。
今回作ったプログラムの実行結果
わりと簡単なプログラムでリアルタイム線分検出ができます。
線分検出カーネルの実装
OpenCLの素晴らしさに感動するには、簡単なカーネルを作ってみるのが一番簡単な方法です。とは言ってもいきなり作るのはハードルが高いので、まずはカメラから画像を取得して線分検出を行うサンプルプログラムを見て、OpenCLの楽しさを分かってもらえると嬉しいです。
プロジェクトの作り方は環境構築編を見てください。実行にはOpenCVが必要なのでNuGetで取ってきてこのサイトを参考に修正すれば動くようになります。
全体的な流れ
OpenCLを使うための鉄板の流れは、次のようになっています。
- プラットフォーム&デバイスの取得
- コンテキスト&コマンドキューの取得
- プログラムのコンパイルとカーネル関数オブジェクトの取得
- メモリ確保&メモリ転送
- カーネル引数の設定&カーネルの実行
デバイスの初期化
OpenCLを使うためにまず最初に行うべきことは、カーネル関数を実行するデバイスを指定し、メモリ取得のためのコンテキストとカーネル実行のためのコマンドキューを取得することです。カーネル関数を実行するプラットフォームとデバイスは、インデックス値で指定します。
使いたいデバイスのインデックス値はC++でOpenCL(環境構築編)で作ったプログラムで調べることができます。今回はオンボードGPU（Intel HD Graphics 4600）を使いたかったので、platformIdxとdeviceIdxに0を指定しています。
// opencl.cpp
// プラットフォーム取得
cl_uint platformNumber = 0;
cl_platform_id platformIds[8];
checkError(clGetPlatformIDs(8, platformIds, &platformNumber));
cl_platform_id platform = platformIds[platformIdx];
// デバイス取得
cl_uint deviceNumber = 0;
cl_device_id deviceIds[8];
checkError(clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 8, deviceIds, &deviceNumber));
gDevice = deviceIds[deviceIdx];
// コンテキスト（メモリ確保などに使用）とコマンドキュー（カーネルの実行などに使用）を作成
gContext = clCreateContext(NULL, 1, &gDevice, NULL, NULL, NULL);
gCommandQueue = clCreateCommandQueue(gContext, gDevice, 0, NULL);
カーネルの初期化
カーネル実行に必要な環境を整えたら、次はカーネルプログラムをコンパイルする必要があります。
ところで余談ですが、OpenCLのカーネルプログラムはテキストデータですし、実行時に好きなタイミングでコンパイルしたり破棄したりできるので、カーネルの動的変更ができそうですよね。機械学習分野とかで使えないのかな...。
// opencl.cpp
// カーネルプログラムのコンパイル
gProgram = compileProgram("kernel.cl");
gKernel = createKernel(gProgram, "detectLine");
コンパイルとカーネル関数オブジェクトの取得をさっくりと行うための関数。
// opencl.cpp
/*
OpenCLのカーネルプログラムをコンパイルして、生成されたプログラムオブジェクトを返します。
*/
cl_program OpenCL::compileProgram(char* fileName)
{
// プログラムの読み込み
FILE* fp;
fp = fopen(fileName, "r");
if (fp == NULL)
{
printf("%s load failed\n", fileName);
return NULL;
}
fseek(fp, 0, SEEK_END);
const int filesize = ftell(fp);
fseek(fp, 0, 0);
char* sourceString = (char*)malloc(filesize);
size_t sourceSize = fread(sourceString, sizeof(char), filesize, fp);
fclose(fp);
// プログラムのコンパイル
cl_program program = clCreateProgramWithSource(gContext, 1, (const char**)&sourceString, (const size_t*)&sourceSize, NULL);
cl_int err = clBuildProgram(program, 1, &gDevice, NULL, NULL, NULL);
// コンパイルに失敗した場合はエラー内容を表示
if (err != CL_SUCCESS)
{
size_t logSize;
clGetProgramBuildInfo(program, gDevice, CL_PROGRAM_BUILD_LOG, 0, NULL, &logSize);
char* buildLog = (char*)malloc((logSize + 1));
clGetProgramBuildInfo(program, gDevice, CL_PROGRAM_BUILD_LOG, logSize, buildLog, NULL);
printf("%s", buildLog);
free(buildLog);
}
free(sourceString);
return program;
}
cl_kernel OpenCL::createKernel(cl_program program, char* kernelName)
{
return clCreateKernel(program, kernelName, NULL);
}
メモリ操作
CPU側とデバイス側ではアドレス空間が異なる（物理メモリが異なる場合もある）ため、CPU側で確保したデータをそのまま使用することはできません。そこで、
- デバイス側にメモリ領域を確保
- CPU側からデバイス側にデータ転送
- カーネル関数でいろいろ計算
- デバイス側からCPU側にデータ転送
- デバイス側のメモリ領域を解放
といった手順が必要になります。まずメモリ領域の確保ですが、もっともシンプルな形で実装すると以下のようになります。デバイス側のメモリにはいろいろ種類があり、高速化のためにはより高度な設定が必要になりますが、とりあえずはこの基本形が大事です。
// opencl.cpp
// デバイスで使用するメモリオブジェクトを確保
gResult = clCreateBuffer(gContext, CL_MEM_READ_WRITE, sizeof(char) * 640 * 480, NULL, NULL);
gOrigin = clCreateBuffer(gContext, CL_MEM_READ_WRITE, sizeof(char) * 640 * 480 * 3, NULL, NULL);
次にclCreateBufferで得られたメモリオブジェクトに対して、データ書き込みおよびデータ読み込みを行います。
// opencl.cpp
clEnqueueWriteBuffer(gCommandQueue, gOrigin, CL_TRUE, 0, sizeof(char) * 640 * 480 * 3, origin, 0, NULL, NULL);
clEnqueueReadBuffer(gCommandQueue, gResult, CL_TRUE, 0, sizeof(char) * 640 * 480, result, 0, NULL, NULL);
最後に、メモリオブジェクトを解放します。
// opencl.cpp
clReleaseMemObject(gResult);
clReleaseMemObject(gOrigin);
カーネル呼び出し
メモリを確保してデータを転送したら、いよいよカーネル関数を呼び出します。カーネル関数を呼び出すのは
- カーネルオブジェクトの引数にメモリオブジェクトを設定
- カーネルの実行サイズを設定
- カーネル実行
といった手順で行います。
// opencl.cpp
cl_int2 size = { 640, 480 };
// メモリオブジェクトをカーネル関数の引数にセット
clSetKernelArg(gKernel, 0, sizeof(cl_mem), &gResult);
clSetKernelArg(gKernel, 1, sizeof(cl_mem), &gOrigin);
clSetKernelArg(gKernel, 2, sizeof(cl_int2), &size);
// カーネルの並列実行数を設定
size_t workSize[2] = { 640, 480 };
// カーネルの呼び出し
clEnqueueNDRangeKernel(gCommandQueue, gKernel, 2, NULL, workSize, NULL, 0, NULL, NULL);
カーネル関数
これが本命のカーネル関数です。今回のプログラムではこのカーネル関数が640×480=307200スレッド起動して、各画素に対して線分検出の計算を行います。線分検出の方法はいたって簡単で、検出対象画素周辺の3×3領域に対して、中心が明るく周辺が暗いほど高い値になるように係数をかけて足し合わせるだけです。明るさが一様な画素では打ち消しあってゼロになりますが、点や線の上にある画素ではゼロより大きな値が得られます。詳しくは網膜の中心周辺拮抗型受容野について調べてください...。
// kernel.cl
__kernel void detectLine(__global unsigned char* result, __global unsigned char* origin, int2 size)
{
// get_global_idは並列実行されたスレッドが何番目のスレッドなのか取得する関数
int x = get_global_id(0);
int y = get_global_id(1);
float val = 0;
val += 6.828 * origin[(size.x * y + x) * 3];
val -= origin[(size.x * y + (x - 1)) * 3];
val -= origin[(size.x * y + (x + 1)) * 3];
val -= origin[(size.x * (y - 1) + x) * 3];
val -= origin[(size.x * (y + 1) + x) * 3];
val -= 0.707 * origin[(size.x * (y - 1) + (x - 1)) * 3];
val -= 0.707 * origin[(size.x * (y - 1) + (x + 1)) * 3];
val -= 0.707 * origin[(size.x * (y + 1) + (x - 1)) * 3];
val -= 0.707 * origin[(size.x * (y + 1) + (x + 1)) * 3];
result[size.x * y + x] = (unsigned char)val;
}
エラーチェック
OpenCLの関数が失敗したとき、戻り値またはエラーを格納する引数にエラーコードが返ってくるので、これを簡単にチェックするためのマクロがあると大変便利です。
// opencl.h
#define checkError(openclFunction) \
if (cl_int err = (openclFunction)) \
{ \
printf("error : %d\n", err); \
}
実行結果
これまでOpenCL独特のもの（デバイス、カーネル、メモリオブジェクトなどなど）がいろいろと出てきましたが、実際のところデバイスの取得やプログラムのコンパイルは定型的な作業なので、慣れてしまえばなんてことはありません。メモリ確保とカーネル実行はとてもシンプルなコードなので、こんな簡単なプログラムでサクサクとした線分検出ができるのは楽しいですよね。
まだ続く
次はこんな素晴らしいOpenCLをさらに手軽に利用すべく、C#からOpenCLを呼び出す方法をやる予定です。さらに、楽しいOpenCLから実用的なOpenCLへ向けて、C++でOpenCL(高速化編)もあります。