今回は謎解明編。
ジッちゃんの名にかけて！
なぜ、raytrace_nocl.mの#if FOR_OPENCL_TRANSLATEのarearender関数が、正しく動く事がOpenCL版で正しく画像が出るかどうかに繋がるのか？
あと、なんでOpenCL未使用版でこの関数を使ったところで全然速くならないのか？
まず、rendering関数で
としてるんですが、こいつで決定してるのが、画像全体の演算をおこなうのにarearender関数を何回呼び出す必要があるかの決定。
にarearender関数で処理する画素数を設定する事で割り出します。
ここではarearender関数呼び出し1回につき、1画素演算で算出してるんで、globalcountは1000x1000て事になるわけですな。
ちなみに、ここではlocalに1を入れているけど、例えば5 x 5の画像をarearender関数で8画素ずつ処理（local = 8）するなら
4回arearender関数を呼び出す事で、すべての画素を演算できることになります。
当然、最後の呼び出しは中途半端な処理数になるわけで、そのため
という式（試しにlocal = 8、width = 5、height = 5で計算機使って計算してみましょう。shellでexpr使ったりしてな、逆にめんどいわ）で、回数を割り出しているわけです。
で、rendering関数では、その回数分arearender関数を呼び出して処理させている。
もちろんarearender関数も、最後の領域が必ずしもlocalで綺麗に割り切れない事を考慮した設計になっとります。そこらへんはサンプルソース見てください。arearender関数先頭で
としてlocalを引数でもらわずにglobalcountから、あらためて計算してるのに特に意味はなく、OpenCLでの高速化が
で考えているので、素直にそれを実装しただけです。ちょっとでもスピード上げたいならlocalそのまま渡した方がいいんですが実験なんでこうしてる。
肝心なのは、次の
で、こっちを引数にしないのは明確な理由があって、OpenCLではこの情報は引数では渡せないんですな。
というのも、これが高速化のキモなわけでOpenCLを使わないraytrace_nocl.mでは単純にrendering関数で
としてループ処理によって順次実行されるarearender関数が
非OpenCL版
OpenCLでは同時に並行して実行されるわけです。
OpenCL版
このため、arearender関数側では、自分がどの領域の画素を受け持つべきかを
で確認するわけです。
このget_global_idってのは本来、OpenCL側が提供する関数であり、呼び出した処理が平行処理用に分割された区画の何番目であるかを教える関数。
raytrace_nocl.mでは、それを自前で模造してるだけなんですよ。
引数で0を渡してるのは、平行処理分割を1次元じゃなく2次元、3次元というように多次元でおこなう事ができるため。例えばint型の2次元配列変数
で、valueの添字が1、2の値と言われた時に、
なのか
なのかわからない。
次元0の添字は1、次元1の添字は2というように、次元を指定して返ってきた値で初めて
とわかるわけで、get_global_idの引数はこの次元の指定にあたります。
今回は平行処理分割を1次元でおこなうので、次元は0を指定するわけです。
注）幾何学上の2次元、3次元と、平行処理の2次元、3次元が本質的に無関係という事に注意せんと駄目っすね。今回なら2次元画像を平行処理してるわけだけど、今回の平行処理自体は1次元でおこなう。
これがraytrace_nocl.mで事前確認した内容。
ちゃんとarearender関数が機能してるのを確認した上で、この処理をOpenCL用のソースに移植します。それがResourcesグループに入れているpixelrender.clというテキストファイル。
開いてもらうとわかるけど、arearender関数と、そこから呼び出される関数がほぼそのまんまコピーされてます。
違いは
というキーワードが付いてる事、#include文が無い事、拡張子が.cでも.mでもなく.clって事ですな。
こいつが何かというと、OpenCL用の言語でして…
C言語をベースにしてるので、enumとか#defineもそのまま使えて、あんまり変更せずに移植できるようになっとるわけです。
ただ～し、Xcodeはこのソースをコンパイルしません。PNG画像同様、アプリケーションバンドルディレクトリのリソースディレクトリにコピーするだけです。
ターゲットグループのデスクロージャを開いてsphere＞バンドルリソースをコピーを開くと見える。
じゃ、誰がコンパイルするのか？
プログラムテキストはコンパイルして演算装置が理解できるマシンコードにしないと、動作してくれないわけで…
そもそも演算装置って何かと考えると、OpenCLでいくぜ！ (2)で説明したようにユーザーの環境でコロコロ変わるわけですよ。
nVIDIAだったりATIだったり、Core i7だったり、Core 2 Duoだったり、Cell Engineだったり、もしくはそのすべてだったり。
つまりは
なわけです。
これがraytrace_cl.mの
というOpenCL用APIなわけですわ。
というわけで次回は、コンパイルしたマシンコードを演算装置で平行で実行させる方法について！
ではでは。
ジッちゃんの名にかけて！
なぜ、raytrace_nocl.mの#if FOR_OPENCL_TRANSLATEのarearender関数が、正しく動く事がOpenCL版で正しく画像が出るかどうかに繋がるのか？
あと、なんでOpenCL未使用版でこの関数を使ったところで全然速くならないのか？
まず、rendering関数で
int local = 1;
int globalcount = (width * height + (local - 1)) / local;
int globalcount = (width * height + (local - 1)) / local;
としてるんですが、こいつで決定してるのが、画像全体の演算をおこなうのにarearender関数を何回呼び出す必要があるかの決定。
local
にarearender関数で処理する画素数を設定する事で割り出します。
ここではarearender関数呼び出し1回につき、1画素演算で算出してるんで、globalcountは1000x1000て事になるわけですな。
ちなみに、ここではlocalに1を入れているけど、例えば5 x 5の画像をarearender関数で8画素ずつ処理（local = 8）するなら
4回arearender関数を呼び出す事で、すべての画素を演算できることになります。
当然、最後の呼び出しは中途半端な処理数になるわけで、そのため
(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関数が
非OpenCL版
OpenCLでは同時に並行して実行されるわけです。
OpenCL版
だから速い！
このため、arearender関数側では、自分がどの領域の画素を受け持つべきかを
get_global_id(0);
で確認するわけです。
このget_global_idってのは本来、OpenCL側が提供する関数であり、呼び出した処理が平行処理用に分割された区画の何番目であるかを教える関数。
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というテキストファイル。
開いてもらうとわかるけど、arearender関数と、そこから呼び出される関数がほぼそのまんまコピーされてます。
違いは
__global
__kernel
__kernel
というキーワードが付いてる事、#include文が無い事、拡張子が.cでも.mでもなく.clって事ですな。
こいつが何かというと、OpenCL用の言語でして…
C言語をベースにしてるので、enumとか#defineもそのまま使えて、あんまり変更せずに移植できるようになっとるわけです。
ただ～し、Xcodeはこのソースをコンパイルしません。PNG画像同様、アプリケーションバンドルディレクトリのリソースディレクトリにコピーするだけです。
ターゲットグループのデスクロージャを開いてsphere＞バンドルリソースをコピーを開くと見える。
じゃ、誰がコンパイルするのか？
プログラムテキストはコンパイルして演算装置が理解できるマシンコードにしないと、動作してくれないわけで…
そもそも演算装置って何かと考えると、OpenCLでいくぜ！ (2)で説明したようにユーザーの環境でコロコロ変わるわけですよ。
nVIDIAだったりATIだったり、Core i7だったり、Core 2 Duoだったり、Cell Engineだったり、もしくはそのすべてだったり。
つまりは
動作する実行環境時に演算装置別にコンパイルが必要
なわけです。
これがraytrace_cl.mの
clGetProgramBuildInfo
というOpenCL用APIなわけですわ。
というわけで次回は、コンパイルしたマシンコードを演算装置で平行で実行させる方法について！
ではでは。