OpenCLでいくぜ! (8)

テーマ:
 仕事立て込んできた上に風邪引きました。
 まじに呪い?

 まあ、いい。

 で、前回言ったような平行で動作する処理を、OpenCLではカーネルって呼んでます。
cl_kernel


$テン*シー*シー-1

 こいつは演算装置の平行動作可能な演算ユニット(たいがいはその下に存在するプロセッシングエレメント)ごとに用意される事になるわけで…

$テン*シー*シー-1

 このカーネルを沢山平行動作させればさせるほど、計算は速くなるわけですわ。
 といっても演算ユニット/プロセッシングエレメントは物理的に限りがあるわけで、いつでも暇してるわけでもないので、その利用状況を把握して、動いてない演算ユニット/プロセッシングエレメントがあればカーネルを割り当て、仕事をさせるという現場監督さんがいるわけですな。

 OpenCLではコマンドキューって呼んでます。
cl_command_queue

 このコマンドキューのおかげで、使う方は「キューちゃ~ん、仕事取っといたから、スケジューリングあとよろ。」って感じで作業を丸投げできるわけです。

$テン*シー*シー-2

 非常に頼もしい現場監督さんなわけですが、それでも最低限の情報は必要なわけで
使える演算装置は何か
使うカーネルは何か
カーネルには何を伝えればいいか
カーネルに伝える方法は

 ここらへんは、教えてあげないとコマンドキューも動きようがない。
 ということでraytrace_cl.mでは以下のようなステップでコマンドキューに仕事を依頼しています。

$テン*シー*シー-3

 コマンドキューは複数の演算装置を管理しないみたいで、Core i7ならCore i7、nVIDIAならnVIDIAのコマンドキューを用意する必要があるみたいっす。
 でもって、演算装置一つをコマンドキュー一つが独占する前提とはしていないので、演算装置の利用状況を把握するマネージャーさんが必要になるみたいでOpenCLではコンテキストって言ってます。
cl_context

 さんざんiPhoneアプリ開発でも出てくるコンテキストですわ。いつも言ってるように「文脈、話の流れ」って解釈しましょう。今、現状は演算装置Aはこれこれに使われてて、演算装置Bは暇してて、というように常に現状を把握するオブジェクトです。
 こっちは複数の演算装置を管理できるようになってて
cl_context clCreateContext(
const cl_context_properties * /* 属性 */,
cl_uint /* 演算装置の数 */,
const cl_device_id * /* 演算装置の配列 */,
void (*pfn_notify)(const char *, const void *, size_t, void *) /* 通知関数 */,
void * /* 通知関数で受け取れる利用者側情報 */,
cl_int * /* エラー情報 */);

 て、ふうに作成します。
 コマンドキュー自体もコンテキストに依頼して作成してもらう関係から、コンテキストはコマンドキューも含め全体の資源の利用状態を把握できるようになってますな。おそらくコマンドキューは内部でコンテキストと連絡し合いながらカーネルの運用をするものと思われ。
cl_command_queue clCreateCommandQueue(
cl_context /* コンテキスト */,
cl_device_id /* 対象の演算装置 */,
cl_command_queue_properties /* 属性 */,
cl_int * /* エラー情報 */);

 コマンドキューを作ったら、次に作るのはカーネル。こいつは前回言ったように演算装置ごとに用意されないといけないので、テキスト(ASCII、UTF8でもOK?)ファイルとして提供されているものを読み込んで対象の演算装置用にコンパイルしてます。
 というわけで、カーネルの作成は以下の数ステップが必要。

$テン*シー*シー-4

 ソースファイルのメモリ上のバッファへの読み込みはファイルパスを特定する以外は、標準的なC言語の読み込み処理。
 コンパイル用ソース(OpneCLではプログラムって呼んでいる)は、複数のバッファを指定できるようになってるけど、今回はpixelrender.clファイルだけなんでバッファの数は1で呼び出しとります。
cl_program clCreateProgramWithSource(
cl_context /* コンテキスト */,
cl_uint /* テキストバッファの数 */,
const char ** /* バッファの配列 */,
const size_t * /* 各バッファのバイト数 */,
cl_int * /* エラー情報 */);

 で、このプログラムを使いたい演算装置を指定してコンパイル。同じく複数の演算装置を指定できるようになってるけど、ここではひとつの演算装置だけ指定してる。あとコンパイラ中の情報をもらう事ができるようになってますが、今回は特になにも指定せず。
cl_int clBuildProgram(
cl_program /* プログラム */,
cl_uint /* 演算装置の数 */,
const cl_device_id * /* 演算装置配列 */,
const char * /* オプション */,
void (*pfn_notify)(cl_program , void * ),/* 通知関数 */,
void * /* pfn_notifyで使える利用者側情報 */);

 ただし、このAPIだけは特別にエラーが返った場合に、どこが間違ってるかを知るためにclGetProgramBuildInfo使ってコンパイルエラーの内容を表示させてます。
if (err != CL_SUCCESS) {
size_t len;
char buffer[2048];
printf("Error: Failed to build program executable!\n");
clGetProgramBuildInfo(program, device_id,
CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
printf("%s\n", buffer);
exit(1);
}

 今のところXcodeはOpenCL側のソースの構文チェックはしてくれないので、こんな風に実行時にチェックするしかないのよ。
 ちなみにフィックスターズという会社から出ているFOXCてのはOpenCLコンパイラで、こいつ使えば事前に構文チェックとかもできそうな…

フィックスターズ

 コンパイルができたらプログラムからカーネル作成。ここではカーネルのエントリーとして使う関数(アプリケーションで言うところのmain関数の役割ですな)の名前を指定して作成です。当然"arearender"ね。
cl_kernel clCreateKernel(
cl_program /* プログラム */,
const char * /* カーネル名称 */,
cl_int * /* エラー情報 */);

 で、ここまででコマンドキュー、カーネルの準備ができたので、残るは
カーネルには何を伝えればいいか
カーネルに伝える方法は

 ということになります。
 伝える情報は当然、
レイトレーシングで作成する画像情報
球体情報
光源情報
視点情報

 となり、これらはカーネルのエントリー関数に渡す引数という形で渡すことになるんだけど、ここで重要な点がひとつ!
使う側のメモリとカーネル側のメモリは必ずしも物理的に同じものではない

 ちゅーことです。
 ま、CPU使う場合、物理的には同じメモリなんだけど(どう扱われるか詳細までは知らない)ビデオカード側のGPUを使うならビデオカード側のメモリにデータがおかれている必要があるわけです。
 そのため以下のように、メモリ間のやり取りが必要になる。

$テン*シー*シー-5

 この演算装置側バッファ(OpenCLではメモリオブジェクトって呼ぶみたい)の確保が
cl_mem clCreateBuffer(
cl_context /* コンテキスト */,
cl_mem_flags /* メモリの属性フラグ */,
size_t /* 大きさ */,
void * /* host_ptr? */,
cl_int * /* エラー情報 */);

 で、そのバッファへの書き込みが
cl_int clEnqueueWriteBuffer(
cl_command_queue /* コマンドキュー */,
cl_mem /* 対象のメモリオブジェクト */,
cl_bool /* 書き込みが完了するまでAPIから返らないかどうか */,
size_t /* データアドレスからのオフセットバイト */,
size_t /* 書き込みバイトサイズ */,
const void * /* データアドレス */,
cl_uint /* 待ちイベントの数 */,
const cl_event * /* 待ちイベント配列 */,
cl_event * /* 自分の処理イベント */);

 てことになるわけですな。
 ここで引き渡している待ちイベントってのは、いわゆるposix threadでいうEventと同じで、コマンドキュー内で先行して平行動作する処理は個別の「待ちイベント」を持っていて、処理が終わった時にイベントという形で通知してくれるようになっている。
 そのため、先行するAとBの処理が終わるまで、自分は動作する事ができないって場合、この「待ちイベント」を使ってAとBの処理が完了するのを待つ事ができるわけです。
 当然clEnqueueWriteBufferも一つの処理なんで、自分の「待ちイベント」を持ってて、必要なら最後の引数で受け取る事もできりるようになってます。
 なんですが、今回は「書き込みが完了するまでclEnqueueWriteBufferから返らない」という指定をしているのと、先行する処理で待つ必要のある処理は無いので、どちらもNULL指定してるわけですね。

 これで準備はOK。
 あとはこれをカーネルのarearender関数に伝える方法ですが、pixelrender.clをあけてもらいarearender関数を見てもらうと、引数のところどころに__globalというキーワードが付いてるのがわかると思います(ちなみにarearenderはカーネルとして利用するという合図の__kernelというキーワードがついている)。
__kernel void arearender(
const int image_width,
const int image_height,
__global float* spheres,
const size_t sphere_count,
__global float* eye_position,
__global float* light,
__global unsigned char* baseAddress,
const size_t bytesPerRow,
const size_t globalcount)

 これが、さっき作っていたメモリオブジェクトに対応してるわけで、このように準備したメモリオブジェクトをカーネルのエントリ関数に引数として渡すための準備がraytrace_cl.mの
cl_int clSetKernelArg(
cl_kernel /* カーネル */,
cl_uint /* 何番目の引数か */,
size_t /* 引数の大きさ */,
const void * /* 引数のアドレス */);

 なわけですわ。

 そしていよいよ実行指示として
cl_int clEnqueueNDRangeKernel(
cl_command_queue /* コマンドキュー */,
cl_kernel /* カーネル */,
cl_uint /* 作業次元 */,
const size_t * /* グルーバル作業オフセット */,
const size_t * /* グルーバル作業数 */,
const size_t * /* ローカル作業数 */,
cl_uint /* 待ちイベント数 */,
const cl_event * /* 待ちイベント配列 */,
cl_event * /* 自分のイベント */);

 を呼び出して実行!となります。
 最後の3つの引数はclEnqueueWriteBufferで説明した待ちイベント関係ね。作業次元も前回説明した処理分割を何次元でやるかって話で、今回は1次元でやってるので1。

テン*シー*シー-2

 1次元で分割するならピンク色の作業を担当してるカーネルでは
get_global_id(0) = 2

 と返るようになってるし

テン*シー*シー-3
 というように2次元で分割するなら
get_global_id(0) = 1
get_global_id(1) = 0

 という形でget_global_idが返ってくるので、これをもとに自分の作業担当を把握するわけだ。

 いまいち使う局面を理解しきれていないのがグルーバル作業数とローカル作業数。
 グルーバル作業数が全体の作業分割数ではあるんですが、ローカル作業数ってのも指定できるようになってるんですわ。
 演算ユニットの中にあるプロセッシングエレメントがローカル作業数なのかな~とか、いろいろ調査中。あと、グルーバル作業数はローカル作業数の整数倍でなければいけないみたい。
 ま、ちょっと謎。
 
 clEnqueueNDRangeKernelを呼び出すと、すぐに返ってくるけど、それはコマンドキューに処理を登録しただけであって、処理が終わったわけではない。細かく制御したいなら待ちイベントをもらって、処理を見張ってもいいんですが、レイトレ終わるまでこっちは何もすることないんで
clFinish

 でコマンドキューの全処理が終わるまで待ってます。この処理が時間かかるようなら、clFinishは使わず、レイトレ演算が終わるのも見張りつつGUIのイベント処理に対応していくって実装が必要でしょうね。

 結果の取り出しには
clEnqueueReadBuffer

 ほとんどclEnqueueWriteBufferと同じなんで説明は割愛。
 そんな感じなり。
AD