OpenCLについてメモ
研究室のMacのGPUがATI Radeonであることもあって、前々からOpenCLに移行しようと思っていたが、nvccのダメダメさ加減に耐えられなくなって、本格的にOpenCLに移行することにした。
カーネルのソースコードをプログラムに埋め込みたい
カーネルのソースコードがファイルとして外部にあるのは、ポータビリティーに欠けるというかなんとなく気持ち悪いので、埋め込んでしまいたいと思った。(逆に、デバッグしたいとかで後から編集したければ埋め込まないほうが良いけど。)
どうやって埋め込むか調べた。hexdumpを使って、ソースファイルをC言語で読み取れるByte配列形式にしちゃおうとまず考えた。それであれこれ調べていたら、xxdというコマンドがあることを知った。
xxd -i < kernel.cl
これを実行すると、kernel.clをunsigned charの配列として吐き出してくれるのだ。この出力を例えばkernel.cl.hとかしておいて
const unsigned char kernel_source[] = { #include "kernel.cl.h" };
とかすれば埋め込める。便利だ。
ちなみにconst char kernel_source[] = …と書くとうまくいかないので要注意。
カーネルをコンパイルしたい。
カーネルのプログラムを見せたくないこともあるかもしれないので、コンパイルされたバイナリを保存して再利用できないか調べた。
自分が試した範囲内では上手く行かなかったのだが、どこかでミスがあったのかもしれないので、メモとして残しておく。
clBuildProgramでカーネルをコンパイルした後に、これを書く。
size_t kernel_bin_size; clGetProgramInfo(program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &kernel_bin_size, NULL); cerr << "Kernel Binary Size:\t" << kernel_bin_size << endl; std::vector<char> kernel_bin(kernel_bin_size); clGetProgramInfo(program, CL_PROGRAM_BINARIES, kernel_bin_size, &(const void*&)kernel_bin.begin().base(), NULL); cout.write(kernel_bin.begin().base(), kernel_bin_size);
kernel_bin_sizeはコンパイルされたカーネルバイナリのサイズを記憶させる。そのサイズ分だけstd::vectorでメモリを確保し(std::vectorなのは単に個人的趣味。mallocでもnewでも構わない)、そこにバイナリを書き出し、最後で標準出力に出力した。別に標準出力じゃなくても構わない。
これでカーネルのバイナリを保存したら、再利用するときはclCreateProgramWithBinaryを使う。
program = clCreateProgramWithBinary(context, 1, &device_id, &source_size, &source_str, &binary_status, &ret);
でも、ちゃんと動かなかった…なんでかな。
ちなみにCUDA環境だとptxらしいものが出力されていた。Appleだとよくわからないバイナリだった。
C++ Bindings
C++ Bindingsされたヘッダーがある。
このページの、cl.hppというのを落としてインクルードすれば良い。
ちなみにMacの場合、/System/Library/Frameworks/OpenCL.framework/Headersにopencl.hなどのヘッダが置いてあり、ここにcl.hppを置いておくのが合理的だと思われる。あと個人的趣味でopencl.hppという名前でcl.hppへのシンボリックリンクをはっておいた。
C++ Bindingsの説明は
にある。単にclという名前空間に押し込んだだけなようだけれども。Open CL v1.1のものなので更新されたら新しいリンクをみましょう。
ちなみにカーネルのソースコードはC++ではなく厳密なC言語で書かないとダメなようだ。例えば、変数宣言は関数の一番最初に書かないとだめ。他にもなにか制約がありそう。例えばconst char data[] = "hello!"; でdataを読み出すのがうまくいかなかったり。
その他
clEnqueueNDRangeKernelで指定する、global_work_sizeは必ずlocal_work_sizeの倍数でないといけない。少なくともAppleでは倍数でないと動作しなかった。
size_t global_work_size = (MEM_SIZE/3)*3; size_t local_work_size = MEM_SIZE/3; ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global_work_size, &local_work_size, 0, NULL, NULL);
global_work_sizeは全work item数を表し、local_work_sizeは各work group内のwork item数を表すものらしい。