カーネルをなんとかして関数テンプレートにしたい

はてなブログに移行しました.そして最初の記事です.

最近 Xeon Phi なるものを触ったりしていますが,かなりとても書きづらいというか pragma ベースで書くのは (OpenMP で楽に並列化できるレベルなら良いですが) 決して楽ではないと僕は感じています.なぜならば pragma ベースで Xeon Phi のコードを実行 (Offload 実行モデル) だとメモリの取り扱いが若干謎で,ここで確保されてるから大丈夫だろうと思ったら SEGV になったりするし,メモリの確保とか解放,コピー操作といったメモリに関する管理をやろうとすると pragma ベースではかなりつらいしめんどくさい,というかそこら中に複雑な pragma が満ちます.それだったらいっその事 OpenCL で書いた方がメモリ管理はこっちでできるし良いと思ったので勉強しています.

幸いにして IntelXeon Phi を動かせる OpenCL SDK を用意していますし,Xeon Phi を使った OpenCL アプリケーションに関するガイドラインもあるので,問題は OpenCL を使った事がない,ということだけです.

ということで若干気に入らないスタイルもあるものの,自分で書くのは面倒なので C++ Wrapper を使いながら,コードを書いているんですが,OpenCL 1.2 以前ではオペレータオーバーロードも,テンプレートもサポートされていないとかいうので辛い*1 *2

私たちは C++ プログラマであるため,同じコードを何度も書くのは全くのナンセンスであり,コピペして型と関数名を変更するなど許されない.それに管理がめんどくさいし,型の数だけ増えればいずれは破綻する.なんとかしてテンプレートみたいな事はしたい.ならばソースファイルで型だけ置換すれば良いではないかという話になるが,それもそれでつらくて困る.

という訳で一時期話題になってその後は追っていない Boost.Compute *3 を調べてみたところ,やはりソースコードを置換しているようで,ただし OpenCL のカーネルビルドオプションで "-DTYPE=float" のように,コンパイル時に型を置換しているっぽい.ただし,これも問題で,TYPE を置換した場合に本来は置換されては困る部分もあるとかそういうこともあるが,カーネルのコンパイル時に分かる事ではあるのでまあ良い事にしましょう.コンパイル万歳.

なので,それを実践するとこうなる.

__kernel
void func(int n, __global TYPE * x) {
  const int tid = get_global_id(0);
  if(tid < n) {
    x[tid] = tid + 1;
  }
}
const std::string kernelSrc = /* std::string でカーネルコードを読み出す */;
cl::Program::Sources source(1, std::make_pair(kernelSrc.c_str(), kernelSrc.size()));
cl::Program program(context, source);
program.build(devices, "-DTYPE=float");
cl::Kernel kernel(program, "func", &err);

これで func は float 型の関数となる.オーバーロードをしたい場合は必要な型の数だけコンパイルして,何らかの方法で (Boost.Compute ではハッシュテーブルか何かで解決していたが) オーバーロード解決してカーネルを呼び出せるようなロジックを作れば良い.

コンパイル時に CPU か CUDA のどちらかをバックエンドにして計算しているコードはあるので,OpenCL のバックエンドを組み立てて上げればパフォーマンスはともかく,そのまま動くはずだ.

*1:OpenCL 2.0 は調べてないので不明.

*2:CUDA ではテクスチャメモリ周りで色々問題が起こったりもするけど使える.CUDA の最新バージョンでは試していない.

*3:OpenCL を用いた Thrust や Intel TBB などと同じ STL-Like なライブラリである.