OpenCLでいくぜ! (9)

テーマ:
 っていってもXcode 4対応ね。

 「OpenCLでいくぜ! (6)」のサンプルのopencl06をXcode 4でビルドしようとするとエラーになるんですわ。

$テン*シー*シー-1

 どうやら
/System/Library/Frameworks/OpenCL.framework/Libraries/openclc

 がないから
pixelrender.cl

 がコンパイルできないよ~ってエラーらしい。
 で、実際どうなのと調べてみたんですがありませんでした。
 調べたい人はFinderの移動>フォルダへ移動…メニューで出るダイアログに
/System/Library/Frameworks/OpenCL.framework/Libraries/

 をコピペして移動ボタンを押すべし。

 本来
pixelrender.cl

 は「OpenCLでいくぜ! (7)」で説明したように実行時にコンパイルされるもんであって、アプリケーション作成時にコンパイルが必要なわけねーんすよ。
 ただ、これまた「OpenCLでいくぜ! (8)」で説明したように事前にソースコードが正しいかを判断する方法が無いって問題もあったわけで…

 考えられるのは、Appleが事前にOpenCL用ソースをコードチェックできるようにOpenCLコンパイラを用意するつもりなのかなって事です。
 で、現段階でOpenCLコンパイラは用意されていないと…
 実際プロジェクトのBuild Rulesには

$テン*シー*シー-3

 OpenCLって項目があって

$テン*シー*シー-2

 とかなってて、これが
/System/Library/Frameworks/OpenCL.framework/Libraries/openclc

 なんじゃないかと思うんですな。
 というわけで.cl拡張子がついてるやつは必ずOpenCLコンパイラが走ることになる。

 今のところ回避方法としてはTARGETSのBuild PhasesタブのCopy Bundle Resourcesを開いてpixelrender.clの項目を選び削除。

$テン*シー*シー-5

 次にAdd Build PhaseでCopy Filesを選んで

$テン*シー*シー-6

 Copy Files Phaseを追加、そこにあらためてpixelrender.clを追加する。

$テン*シー*シー-7

$テン*シー*シー-8

 そうすることでpixelrender.clは単なるファイルのバンドルへのコピーとなって、OpenCLコンパイラが走らなくなる。
 NOTE:つまり、Copy Bundle Resourcesは単にコピーじゃなく、拡張子に基づいたコンパイル作業をともなうコピー処理なわけですな(Build RulesのInterface Builder XIB Compiler、Interface Builder NIB PostProcessorってのがある事に注目)。

 まあ、とにかくこれでめでたくXcode4でサンプルが動くようになるわけです。
 ついでに光源を動かしてアニメーションさせてみた。

------------
サンプルプロジェクト:opencl07.zip
AD

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

OpenCLでいくぜ! (7)

テーマ:
 今回は謎解明編。
 ジッちゃんの名にかけて!

 なぜ、raytrace_nocl.mの#if FOR_OPENCL_TRANSLATEのarearender関数が、正しく動く事がOpenCL版で正しく画像が出るかどうかに繋がるのか?
 あと、なんでOpenCL未使用版でこの関数を使ったところで全然速くならないのか?
 まず、rendering関数で
int local = 1;
int globalcount = (width * height + (local - 1)) / local;

 としてるんですが、こいつで決定してるのが、画像全体の演算をおこなうのにarearender関数を何回呼び出す必要があるかの決定。
local

 にarearender関数で処理する画素数を設定する事で割り出します。
 ここではarearender関数呼び出し1回につき、1画素演算で算出してるんで、globalcountは1000x1000て事になるわけですな。
 ちなみに、ここではlocalに1を入れているけど、例えば5 x 5の画像をarearender関数で8画素ずつ処理(local = 8)するなら

$テン*シー*シー-1

 4回arearender関数を呼び出す事で、すべての画素を演算できることになります。

$テン*シー*シー-2

 当然、最後の呼び出しは中途半端な処理数になるわけで、そのため
(width * height + (local - 1)) / local;

 という式(試しにlocal = 8、width = 5、height = 5で計算機使って計算してみましょう。shellでexpr使ったりしてな、逆にめんどいわ)で、回数を割り出しているわけです。
  で、rendering関数では、その回数分arearender関数を呼び出して処理させている。

 もちろんarearender関数も、最後の領域が必ずしもlocalで綺麗に割り切れない事を考慮した設計になっとります。そこらへんはサンプルソース見てください。arearender関数先頭で
size_t localsize = (image_width * image_height) / globalcount;

 としてlocalを引数でもらわずにglobalcountから、あらためて計算してるのに特に意味はなく、OpenCLでの高速化が
全画像を何分割するか

 で考えているので、素直にそれを実装しただけです。ちょっとでもスピード上げたいならlocalそのまま渡した方がいいんですが実験なんでこうしてる。
 肝心なのは、次の
get_global_id(0);

 で、こっちを引数にしないのは明確な理由があって、OpenCLではこの情報は引数では渡せないんですな。
 というのも、これが高速化のキモなわけでOpenCLを使わないraytrace_nocl.mでは単純にrendering関数で
for (global_id = 0; global_id < globalcount; global_id++)

 としてループ処理によって順次実行されるarearender関数が

テン*シー*シー-3
非OpenCL版

 OpenCLでは同時に並行して実行されるわけです。

テン*シー*シー-4
OpenCL版
 だから速い!

 このため、arearender関数側では、自分がどの領域の画素を受け持つべきかを
get_global_id(0);

 で確認するわけです。
 このget_global_idってのは本来、OpenCL側が提供する関数であり、呼び出した処理が平行処理用に分割された区画の何番目であるかを教える関数。

$テン*シー*シー-5

 raytrace_nocl.mでは、それを自前で模造してるだけなんですよ。
 だから速くならない!


 引数で0を渡してるのは、平行処理分割を1次元じゃなく2次元、3次元というように多次元でおこなう事ができるため。例えばint型の2次元配列変数
int value[3][3];

で、valueの添字が1、2の値と言われた時に、
value[1][2]

なのか
value[2][1]

なのかわからない。
 次元0の添字は1、次元1の添字は2というように、次元を指定して返ってきた値で初めて
value[2][1]

 とわかるわけで、get_global_idの引数はこの次元の指定にあたります。
 今回は平行処理分割を1次元でおこなうので、次元は0を指定するわけです。

 注)幾何学上の2次元、3次元と、平行処理の2次元、3次元が本質的に無関係という事に注意せんと駄目っすね。今回なら2次元画像を平行処理してるわけだけど、今回の平行処理自体は1次元でおこなう。

 これがraytrace_nocl.mで事前確認した内容。
 ちゃんとarearender関数が機能してるのを確認した上で、この処理をOpenCL用のソースに移植します。それがResourcesグループに入れているpixelrender.clというテキストファイル。

$テン*シー*シー-6

 開いてもらうとわかるけど、arearender関数と、そこから呼び出される関数がほぼそのまんまコピーされてます。
 違いは
__global
__kernel

 というキーワードが付いてる事、#include文が無い事、拡張子が.cでも.mでもなく.clって事ですな。
 こいつが何かというと、OpenCL用の言語でして…
 C言語をベースにしてるので、enumとか#defineもそのまま使えて、あんまり変更せずに移植できるようになっとるわけです。

 ただ~し、Xcodeはこのソースをコンパイルしません。PNG画像同様、アプリケーションバンドルディレクトリのリソースディレクトリにコピーするだけです。

$テン*シー*シー-7

 ターゲットグループのデスクロージャを開いてsphere>バンドルリソースをコピーを開くと見える。

 じゃ、誰がコンパイルするのか?
 プログラムテキストはコンパイルして演算装置が理解できるマシンコードにしないと、動作してくれないわけで…
 そもそも演算装置って何かと考えると、OpenCLでいくぜ! (2)で説明したようにユーザーの環境でコロコロ変わるわけですよ。
 nVIDIAだったりATIだったり、Core i7だったり、Core 2 Duoだったり、Cell Engineだったり、もしくはそのすべてだったり。
 つまりは 
 動作する実行環境時に演算装置別にコンパイルが必要

 なわけです。
 これがraytrace_cl.mの
 clGetProgramBuildInfo

 というOpenCL用APIなわけですわ。

 というわけで次回は、コンパイルしたマシンコードを演算装置で平行で実行させる方法について!
 ではでは。
AD