Boost.Computeでグラボを燃やす

はじめに

この記事ではBoost.Computeの紹介と使い方の説明をします。

Boost.Compute

Boost.Computeとは

kylelutz氏が作成しているOpenCLのC++ラッパです。Boostにはまだ正式に採用されていませんが、主要な部分の実装は既に終わっており、十分使えるレベルに達しています。

とりあえず使用例

#include <vector>
#include <algorithm>

#include <boost/compute/algorithm/transform.hpp>
#include <boost/compute/container/vector.hpp>
#include <boost/compute/functional/math.hpp>

namespace compute = boost::compute;

int main()
{
    // デフォルトのdeviceを取得して、contextとcommand_queueを作成
    compute::device device = compute::system::default_device();
    compute::context context(device);
    compute::command_queue queue(context, device);

    // hostでランダムな値を持つvectorを作成
    std::vector<float> host_vector(10000);
    std::generate(host_vector.begin(), host_vector.end(), rand);

    // deviceに領域を確保
    compute::vector<float> device_vector(host_vector.size(), context);

    // hostからdeviceにデータをコピー
    compute::copy(
        host_vector.begin(), host_vector.end(),
        device_vector.begin(),
        queue
    );

    // 各要素の平方根をdevice上で並列に求める
    compute::transform(
        device_vector.begin(), device_vector.end(),
        device_vector.begin(), // in-place
        compute::sqrt<float>(),
        queue
    );

    // deviceからhostへデータをコピー
    compute::copy(
        device_vector.begin(), device_vector.end(),
        host_vector.begin(),
        queue
    );
}

返り値を確認するようなエラー処理やリソース管理をする必要がないため、そのままOpenCLを使うよりはるかに簡単です。

この例ではtransformを使いましたが、その他のSTLアルゴリズムに対応するアルゴリズム（accumulate、sort、find、max_elementなど）も用意されています。

OpenCLカーネルのカスタマイズ

用意されたアルゴリズムを組み合わせるだけでも様々な計算ができますが、アルゴリズムにカーネルを指定する必要がある場合や、アルゴリズムのデフォルトの挙動を変えたい場合、自分で書いたカーネルを直接使いたい場合もあるでしょう。

そうした場合のためにBoost.Computeではカーネルをカスタマイズする方法が複数用意されています。

ラムダ式を使う

ここで言うラムダ式とはC++11で言語機能として導入されたものではなく、Boost.Phoenixなどで用いられている式テンプレートによるものを指します。

次のように使います。

using boost::compute::lambda::_1;
using boost::compute::lambda::_2;
using boost::compute::lambda::sqrt;
using boost::compute::lambda::exp;

boost::compute::transform(
    input1.begin(), input1.end(),
    input2.begin(),
    input1.begin(),
    sqrt(_1)*_2+42 // ラムダ式
);
auto num = boost::compute::count_if(
    input1.begin(), input1.end(),
    exp(_1) <= 666 // ラムダ式
);

余談ですが、内部実装としては式テンプレートからOpenCLカーネルコードを動的に生成しています。

アルゴリズム用のカーネルを書く

ちょっとしたカーネルなら簡単に書くこともできます。

BOOST_COMPUTE_FUNCTION(int, add_four, (int x),
{
    return x + 4; // OpenCLコード
});

boost::compute::transform(
    input.begin(), input.end(),
    output.begin(),
    add_four
);

カーネルを書いてcommand_queueで実行する

より複雑なカーネルを実行したい場合は、OpenCLの作法に従って自分で書いたカーネルを実行することもできます。

const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
    // OpenCLカーネルコード
    __kernel void foo(int k, __global int *x, __global int *y)
    {
        x[get_global_id(0)] = -k*y[get_global_id(0)];
    }
);

auto foo_kernel =
    boost::compute::kernel::create_with_source(source, "foo", context);

boost::compute::vector<int> x(16, context);
boost::compute::vector<int> y(16, context);

... // xの初期化（省略）

foo_kernel.set_args(42, x, y);

queue.enqueue_task(foo);

この例ではcommand_queue::enqueue_taskを使いましたが、多次元データに対してはcommand_queue::enqueue_nd_range_kernelなどもあります。詳しくはBoost.Computeのドキュメントを参照してください。

使用上の注意点

hostからdeviceメモリに直接アクセスしない

boost::compute::vectorなどにはoperator[]演算子が定義されており、添字を指定して要素に直接アクセスすることができますが、時間がかかる上に自分で同期を取らなければならないため止めましょう。

代わりに、copyアルゴリズムなどで範囲ごとに一遍にアクセスするようにします。

In-order実行する

command_queueはデフォルトではIn-order実行するように指定されます。In-order実行であれば同じcommand_queueに入れられたカーネルは入れられた順番に実行され、前のカーネルが終了した後に後のカーネルが実行されます。そのため、後に入れたカーネルが前に入れたカーネルの実行結果に依存していてもデータ競合を起こしません。

一方、command_queueのオプションでOut-of-order実行を指定するとカーネルの実行がcommand_queueに入れられた順番と関係なく実行されるようになります。これは高度な計算の最適化のためには必要なことですが、プログラムの設計が著しく難しくなります。

簡単のために、なるべくIn-order実行することを強くおすすめします。以降、この記事ではIn-order実行を前提にします。

In-order実行で複数のカーネルを並列に実行したい場合は、それ以前の処理が終わるのをwait()やfinish()で待ってから、カーネルごとにqueueを用意してそれぞれで実行し、wait()やfinish()で並列に実行している全てのカーネルの処理が終わるのを同期します。

// kernelAは依存なし
// kernelB1はkernelAの結果に依存
// kernelB2はkernelAの結果に依存
// kernelCはkernelB1とkernelB2の結果に依存
// 各カーネルに引数は既にセットされていると仮定する
auto queue = boost::compute::system::default_queue();
auto queue2(queue.get_context(), queue.get_device());

auto kernelA_event = queue.enqueue_nd_range_kernel(
    kernelA, .../*略*/...
);

// kernelB1とkernelB2は並列に実行できる。
// ここでkernelAが終了するのを待つ必要はない。同じqueueなのでIn-order実行される。
auto kernelB1_event = queue.enqueue_nd_range_kernel(
    kerlenB1, .../*略*/...
);
kernelA_event.wait(); // ここでkernelAが終了するのを待つ必要がある。
auto kernelB2_event = queue2.enqueue_nd_range_kernel(
    kernelB2, .../*略*/...
);

kernelB1_event.wait(); // kernelB1が終了するのを待つ
kernelB2_event.wait(); // kernelB2が終了するのを待つ
queue.enqueue_nd_range_kernel(kernelC, .../*略*/...);

一方で、OpenCLコードのコンパイラを信頼して、全部順番に処理するように書くのも手です。

もしコンパイラが十分賢ければ、kernelB1とkernelB2には依存性がないことを見ぬいて、並列に効率よく実行するようにしてくれるかもしれません（保証はありませんが）。

// kernelAは依存なし
// kernelB1はkernelAの結果に依存
// kernelB2はkernelAの結果に依存
// kernelCはkernelB1とkernelB2の結果に依存
// 各カーネルに引数は既にセットされていると仮定する
auto queue = boost::compute::system::default_queue();

queue.enqueue_nd_range_kernel(kernelA, .../*略*/...);

// kernelB1とkernelB2は並列に実行できるが、あえてそうしない。
queue.enqueue_nd_range_kernel(kerlenB1, .../*略*/...);
queue.enqueue_nd_range_kernel(kernelB2, .../*略*/...);

queue.enqueue_nd_range_kernel(kernelC, .../*略*/...);

最初の実装ではこのように全部順番に処理するように書くことをおすすめします。シンプルで分かりやすく、ミスしにくいためです。

後で性能に問題があれば、深刻な部分から順に明示的に並列化すればいいのです。 ただし、明示的に並列化することでコンパイラによる最適化が阻害される恐れもあるため、明示的に並列化した後は再度計測を行うべきでしょう。

コピーは特別

アルゴリズムの中でもコピーは特別です。hostからdevice、deviceからhostへのコピーはhost-device間で同期されるからです（逆に言うと他のアルゴリズムは基本的に非同期に実行される（ただし、コピーと同じく条件によっては同期実行されるものもある。accumulateとか））。

// 同期バージョンのコピー。コピーが終わるまで処理が進まない。
boost::compute::copy(
    host_vec.begin(), host_vec.end(),
    device_vec.begin(),
    queue
);
... // 何かhost上での処理（前にあるコピーが終了するまで実行されない）

データ量やメモリ帯域の制限によってコピーには時間がかかることがあるため、非同期にコピーしたい場合があります。

そのような場合にはcopy_asyncを使います。

// 非同期バージョンのコピー。コピーの終了を待たずに直ちに処理が返ってくる。
auto copy_event = boost::compute::copy_async(
    host_vec.begin(), host_vec.end(),
    device_vec.begin(),
    queue
);
... // 何かhost上での処理（この裏ではdeviceへのコピーが実行されている……）
... // host_vecにアクセスしないように注意。
copy_event.wait(); // eventをwait()して同期する。
// これ以降はhost_vecにアクセスしても大丈夫

なるべく明示的にcommand_queueを指定する

用意されているアルゴリズムでは実はcommand_queueを指定する必要は必ずしもありません。指定しなかった場合はboost::compute::system::default_queue()が代わりに指定されます。

また、boost::compute::vectorの要素数を指定するコンストラクタもcontextを指定しなかった場合は、boost::compute::system::default_context()が自動で用いられるため省略できます。

そのため、冒頭のコード例は次のように書けます。

#include <vector>
#include <algorithm>

#include <boost/compute/algorithm/transform.hpp>
#include <boost/compute/container/vector.hpp>
#include <boost/compute/functional/math.hpp>

namespace compute = boost::compute;

int main()
{
    // hostでランダムな値を持つvectorを作成
    std::vector<float> host_vector(10000);
    std::generate(host_vector.begin(), host_vector.end(), rand);

    // deviceに領域を確保
    compute::vector<float> device_vector(host_vector.size());

    // hostからdeviceにデータをコピー
    compute::copy(
        host_vector.begin(), host_vector.end(),
        device_vector.begin()
    );

    // 各要素の平方根をdevice上で並列に求める
    compute::transform(
        device_vector.begin(), device_vector.end(),
        device_vector.begin(), // in-place
        compute::sqrt<float>()
    );

    // deviceからhostへデータをコピー
    compute::copy(
        device_vector.begin(), device_vector.end(),
        host_vector.begin()
    );
}

こちらのほうがシンプルで分かりやすいですが、後で明示的に並列実行させたくなった場合や複数デバイスに対応させようとした場合に修正しにくいです。

また、後述するboost::compute::vectorのコピーの話もあるため、なるべく明示的にcommand_queueを指定することをおすすめします。

挙動がおかしい場合

OpenCLのプログラムを書いていて、同じパラメータを与えているはずなのに計算結果が毎回変わったり、エラーが出てプログラムの実行に失敗する場合は、データ競合を起こしている可能性が高いです。

ただし、同じプログラムでも異なるdeviceで走らせた場合（例えば、GPUとCPUで走らせた場合）に浮動小数点の扱いが異なったりするため、計算結果が異なることはよくあります。これはバグではありません。

OpenCLコードのデバッグは通常困難ですが、いくつか僕が採用している指針を示します。

カーネルがデータ競合していないかチェックする

自分で書いたカーネルそれ自体がデータ競合を引き起こしているかもしれません。

特にカーネルのソースコードのメモリに書き込んでいる部分の添字を見て、デバイス上の異なるプロセッサが同時に同じメモリにアクセスしていないかチェックしましょう。

異なるcommand_queueに順番に実行したい処理を入れてしまっていないかチェックする

先程も説明したとおり、異なるcommand_queueに入れられたカーネルが実行される順番は定まりません（deviceがひとつしかなくてもdevice上のプロセッサは複数あるため並列に実行される可能性がある）。

順番に実行したいなら同じcommand_queueに入れるか、明示的にeventをwait()するか、command_queueをfinish()しましょう。

特に、boost::compute::vectorをコピーしようとして次のようなコードを書いてしまうとアウトです。

// 悪い例
some_kernel.set_arg(device_vec1);
boost::compute::system::default_queue().enqueue_nd_range_kernel(
    some_kernel, .../*省略*/...
);
auto device_vec2(device_vec1); // データ競合が発生

boost::compute::vectorのコピーコンストラクタはboost::compute::system::default_queue()ではなくインスタンスごとに独立して内部に持っているcommand_queueを使用するため、カーネルの実行が終わる前にデータがコピーされてしまいます。

// 良くない例
some_kernel.set_arg(device_vec1);
boost::compute::system::default_queue().enqueue_nd_range_kernel(
    some_kernel, .../*省略*/...
);
auto device_vec2(device_vec1.begin(), device_vec1.end()); // OK

今度はboost::compute::system::default_queue()が使用されるため、問題ありません。

しかし、どのようにすればdefault_queue()が使われるのか、いつ独自のcommand_queueが使われるのかなどということをいちいち気にしていたら埒が明きません。

やはり、先程も説明したとおり明示的にcommand_queueを指定するようにしましょう。

// 良い例（おすすめ）
some_kernel.set_arg(device_vec1);
boost::compute::system::default_queue().enqueue_nd_range_kernel(
    some_kernel, .../*省略*/...
);
auto device_vec2(
    device_vec1.begin(), device_vec1.end(),
    boost::compute::system::default_queue() // command_queueを明示的に指定。もちろんOK
);

もしくは明示的に同期を取りましょう。

// 良い例
some_kernel.set_arg(device_vec1);
boost::compute::system::default_queue().enqueue_nd_range_kernel(
    some_kernel, .../*省略*/...
);
boost::compute::system::default_queue().finish(); // 明示的に同期
auto device_vec2(device_vec1); // データ競合は発生しない。

前者の明示的にcommand_queueを指定する方法をおすすめします。うっかり忘れた場合でも間違いを発見しやすく、効率的に実行される可能性がより高いためです。

同じプログラムを複数回実行して値が変わらないかチェックする

データ競合が起きている場合は、計算結果が実行するたびに異なることが多いです。そのため、複数回実行することでデータ競合を検出できます。

ただし、deviceによってはデータ競合があっても見た目上はきちんと動作する場合があります。このことを考えれば、複数の異なるdeviceでチェックするのが望ましいでしょう。

問題のある部分を切り分ける

これはデバッグする上での基本的な手法ですが、使用できるツールが限られているため、OpenCLプログラミングでは特に重要になります。

おわりに

今回は書きませんでしたが、Boost.Computeには他にもOpenGLやOpenCVと組み合わせるための機能や任意のユーザ型をOpenCLでデータ型として扱えるようにする機能、float4などのベクトルデータ型を扱う機能などもあります。

Kylelutz氏のブログやBoost.Computeのexampleに実際のそうしたコード例があるので、参考にしてみてください。

Boost.Computeを使えばGPUを含むOpenCLに対応したあらゆるデバイスを簡単に燃やすことができるので、大量のデータを一律に処理したい場合に使ってみては如何でしょうか。