OpenCLによるGPGPU入門③_二次元配列による計算。

DirectShowで取得した画像をOpenCLで画像処理するプログラムを書いた。
（……といっても、まだ、カラーを白黒に変換するだけのテストプログラム。
GPGPU入門といっても、DirectShowを使っている部分は全く別の話（記事）になるのだが、
global_work_sizeに限度がある事に気が付かずに手間取ったついでに、
OpenCLでの二次元配列の扱いについてメモしておく。）

始めは画像を一次元配列として扱い「global_work_size」を「0～1980*1080-1」で回そうとしたのだが、
global_work_sizeが大きすぎるとエラーを吐くので、結局二次元配列として扱い処理した。
（基本的にはglobal_work_sizeに制限は無いのだが、程度問題、という事で。。。）


以下OpenCLでの二次元配列の扱い方＆ソースファイル。



二次元で計算させるには、「clEnqueueNDRangeKernel()」関数において、
const cl_uint work_dim=2; const size_t global_work_offset[2]={0,0}; //とりあえず「global_work_offset」と「global_work_size」の大小関係に関わらず、「clEnqueueNDRangeKernel()関数」はエラーを返さない。 const size_t global_work_size[2] ={1920,1080}; //{Bitmap_Size};//タスクの総数 //global_work_sizeはlocal_work_sizeで割り切れる必要がある。(この時、global_work_offsetの値は影響しない。) const size_t local_work_size[2] ={1,1}; //1グループあたりのタスクの数 //第6引数に直接0を指定した場合は、「global_work_size」が「local_work_size」で割り切れずともエラーを吐かないが、この配列を通して指定すると、エラーを吐く。 //「local_work_size」の指定が問題となるのは、ワークグループを分割した場合だが、おそらくワークグループを分割する事によってローカルIDを使えるようになる事に、分割のメリットがある。ただ、インデックス空間の次元が増える事との違いが今一ピンと来ていないのでいいかげんな話だなぁ……。 //恐らくだけど、カーネル関数の部分で、「get_local_id()関数」とかを使うんじゃないかな?その時「get_work_dim()」を使って(有効なのは0～(get_work_dim()-1)[次元])次元を所得するとかしないとか? ErrorCode = clEnqueueNDRangeKernel( O_CL_PR->gpu_command_queue, //命令を発行するコマンドキューを指定 O_CL_PR->my_kernel, //ビルドされたプログラムオブジェクトがセットされた、カーネルオブジェクトを指定 work_dim, //カーネルを実行するインデックス空間の次元 global_work_offset,//NULL, //ワークアイテムのオフセットを配列として指定する。　　　 (この値<から>、OpenCL C言語部分において、「get_global_id()関数」よりidが発行される。) global_work_size, //　　　ワークアイテムの総数を配列として指定する。(並列化の数?)(この値<まで>、OpenCL C言語部分において、「get_global_id()関数」よりidが発行される。) //インデックス空間の各次元におけるワークアイテムの数を配列として指定する。 local_work_size,//0, //ワークグループを分割する場合に、一つのワークグループに含まれるワークアイテムの数を指定する。??? 0, //event_wait_list(すぐ下の引数)に指定されるイベントオブジェクトの数 NULL, //ここに指定したイベントオブジェクトが終了するまで、コマンドの実行を待機する。 NULL //カーネルを実行した際に、ここに指定したイベントオブジェクトに対してイベントを発生させる(通知する?)。 );

要するに、
work_dim=2;
で次元を二次元にして、
global_work_offset  と
global_work_size     と
local_work_size      を、
それぞれ要素数2の配列として渡してやればいい。


kernelの方のプログラムは、結局一次元配列として扱う事にしたので、
__kernel void addVector( __global uchar *Bitmap) { size_t index_x = 3*( get_global_id(0) + 1920*get_global_id(1) ); //配列が、N==1,2,3次元の時に使う。 //グローバルID、時にはローカルIDも……使える？、によって計算をコントロールする。 //size_t index_y = get_global_id(1); //配列が、N== 2,3次元の時に使う。 //グローバルID、時にはローカルIDも……使える？、によって計算をコントロールする。 //size_t index_z = get_global_id(2); //配列が、N== 3次元の時に使う。 //グローバルID、時にはローカルIDも……使える？、によって計算をコントロールする。 //グレースケール化 unsigned char grayscale = ( Bitmap[ 0 + index_x ] + Bitmap[ 1 + index_x ] + Bitmap[ 2 + index_x ] ) / 3; Bitmap[ 0 + index_x ] = grayscale; Bitmap[ 1 + index_x ] = grayscale; Bitmap[ 2 + index_x ] = grayscale; }

get_global_id(0) + 1920*get_global_id(1)
として回せばいい。
三倍しているのはRGBが三色で、一回の操作で三色ずつデータ処理しているため。

一応ソースでも上げておくか……。
MD5 | 4e98311256c1255f05aeddfb1374c078 | 12.7 MB | OpenCLによるUBSカメラのリアルタイム画像処理。
http://ux.getuploader.com/ADMIS/download/48/Image_Processing_Ver03.01_00_OpenCL.zip