constexprでOpenCL
最近書いていたコードをgithubに上げてみた。
初githubというか、自分で書いたコードをネットに晒すのも、これが初かも。
動機
OpenCLはデバイス側で動かすコードを、文字列でAPIに渡す訳ですが、これをconstexpr計算（というかsprout::string）で作ってみたらどうかというテスト。
あたりのモチベーションで、色々調べたり書いたりしていたのですが。まさか、とりあえず最初から最後まで処理が繋がるまでに、半年も消費する事になるとは。
できること
やっている事は、型システムを内外で繋げる事がほとんどすべてです。
あとはpolicyベースでバッファ系関数の引数フラグを決めるとか（現在スタブ）、新標準ライブラリのスレッド制御によるロックの隠蔽といった辺りを少しやってみた。
こんな感じでカーネルを宣言し(本当はこいつらにconstexprを付けないとconstexpr計算にならないのですが、付けると何故かgccが内部エラーする)、
auto fill_index = code::func(
"fill_index",
code::returns<pfm::void_>()|
code::param<code::global<pfm::int_>*>("pInt"),
"int id = get_global_id(0);\n"
"pInt[id] = id;"
);
auto twice = code::func(
"twice",
code::returns<pfm::void_>()|
code::param<code::global<pfm::int_>*>("pInt"),
"int id = get_global_id(0);\n"
"pInt[id] = pInt[id] * 2;"
);こんな感じで呼び出し、
queue(run_kernel(program, fill_index(bufWrite), item_count));
queue(run_kernel(program, twice(bufWrite), item_count));こんな感じでCPU側から結果を取り出せます。
auto future2 =
queue(bufConst.with_range(
[](iterator i_begin, iterator i_end){
return std::accumulate(i_begin, i_end, 0);
}));
chrono::steady_clock::time_point tp2 = chrono::steady_clock::now();
std::future_status result2 = future2.wait_until(tp2 + chrono::seconds(5));
assert(result2 == std::future_status::ready);
assert(future2.get() == item_count * (item_count - 1) / 2 * 2);(いずれもtest/buffer.cpp内のコード)
別案
別案として、アルゴリズム自体をexpression templateで記述するような道もあって、調べてみるとアイデアがちょこちょこ引っ掛かりました*2。
しかし、OpenCLの中身を見てみると、メモリバリア等のローレベルな部品が色々あり、処理の本体はあくまで順番通りに書けるようにしておく必要がありそうで、結果こうしてみました。
