読者です 読者をやめる 読者になる 読者になる

Project German-style

同人、Github、その他活動の記録と報告

constexprでOpenCL

C++

最近書いていたコードをgithubに上げてみた。

初githubというか、自分で書いたコードをネットに晒すのも、これが初かも。

動機

OpenCLはデバイス側で動かすコードを、文字列でAPIに渡す訳ですが、これをconstexpr計算(というかsprout::string)で作ってみたらどうかというテスト。

  • せっかくFusion*1 APUを買ったのでGPU動かしてみたい
  • C++11使ってみたい

あたりのモチベーションで、色々調べたり書いたりしていたのですが。まさか、とりあえず最初から最後まで処理が繋がるまでに、半年も消費する事になるとは。

できること

やっている事は、型システムを内外で繋げる事がほとんどすべてです。

あとは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の中身を見てみると、メモリバリア等のローレベルな部品が色々あり、処理の本体はあくまで順番通りに書けるようにしておく必要がありそうで、結果こうしてみました。

今後

今後はとりあえずAPIの網羅を目指すか、サンプルコードと称して面白いGPGPUアルゴリズムを研究するか、隠れた野望である高階関数を実装するか、といった方向性で、まだまだ弄って行きたいです。

*1:若干キラキラネームなライブラリ名は、Fusion APUとかFusion System Architecture(Heterogeneous System Architectureの旧称)に由来しているらしい

*2:Haskellの例ですが[http://amddevcentral.com/afds//assets/presentations/2910_3_final.pdf:title=こんなの]とか