SSブログ

OS XのOpenCL - その24 [OS XのOpenCL]

これまでやってきて、特にカメラ画像に対する処理を自分でも書いてみて、OpenCLって結構難しいということがわかってきた。コンパイラが最適化するCPU用のコードに比較して十分なメリットを出すのがすごく難しい。原因を知るために自分の書いたコードを詳しく見てみるとどうもメモリアクセスが律速してるとしか思えない。CPUの場合はメモリアクセスのバンド幅が広いのと、最近のコンパイラが間抜けなコードを書き直してくれるのでよほど変な書き方をしない限りほぼ最高のパフォーマンスが得られる(ただしもちろん、間抜けなアルゴリズムまでは書き直してくれない)。確かに、こないだ訳し終わったガイドにあったガウシアンブラーではメモリアクセスの最適化が一番効く、という結論にも読める。このガウシアンブラーのチューニングをもっと詳しく調べて自分の書くコードに応用できるようにしないといけない。

ということでOS XのOpenCLはそろそろおしまいにしようと思うんだけど、ガイドを読んでて面白いと思った分岐処理。ガイドではさらっと書いてあるけど、僕なりに理解したのを書いて未来の自分用に残しておく。

15.3  カーネル内の分岐命令

カーネルのコードには分岐、ようするにif文が含まれていてもいい、ということになっている。ハードウェア的には条件判断もビット操作に展開されるので、分岐を禁止するとプログラミング不能になる、というわけではないが、極端に面倒ないわゆるボイラープレートコードが大量に必要になってしまう。そのくらいならGPU側で細工して分岐命令を実行できるようにしよう、ということだろう。

しかし、よく考えると、例えばバイナリツリーの一番下の葉の部分を並列に処理するようなカーネルを考えた場合、そしてもしその葉が右と左のどちらについているか、さらにその親がどちらについているかで異なる処理が必要になる場合、葉は全て違う処理をすることになる。これを並列に処理するにはSIMDではなくMIMDになるということである。それを効率よく処理できるできるようなGPUはないらしい。

GPUがSIMDマシンでありながら分岐命令を含めることができるというのはどういうことかというと、例えば
    int     x = get_global_id(0);
    if (x == 0)
        a = 1;
    else
        a = -1;
なんていうコードがカーネルにあったとする。条件を満足するエレメントと満足しないエレメントでは実行している命令が違う、ということになってしまう。そのかわり
    int     x = get_global_id(0);
        a = 1;   // x == 0 の場合aへの書き込みを禁止する
        a = -1;  // x != 0 の場合aへの書き込みを禁止する
というようなことをしているらしい。全てのエレメントは同じコードを実行して、条件分岐は別のところで制御する、ということのようである(すべてのGPUがそのようなやり方をしているかどうかはわからない)。

従って例えば
    int     x = get_global_id(0);
    if (x == 0)
        proc0();
    else if (x == 1)
        proc1();
    else if (x == 2)
        proc2();
    ....
なんていうコードはSIMDではなくすべてのエレメントでproc0()、proc1()、proc2()...が順番に実行されることになって、実際にはまったく並列化されない。

つまり、カーネルのコードは分岐の書き方で効率が何倍も変わることがありえる、ということだろう。たとえば
    int     x = get_global_id(0);
    if (x == 0)
        func(a);
    else
        func(b);
というコードがあったとする。これは
    int     x = get_global_id(0);
    int     arg
    if (x == 0)
        arg = a;
    else
        arg = b;
    func(arg);
としたほうが効率はいい、ということになる。もしfunc()呼び出しが重いならパフォーマンスは大きく違ってくる。

面白いけどうっかり気にせず書きそうで、注意しないといけない。
nice!(0)  コメント(0)  トラックバック(0) 

nice! 0

コメント 0

コメントを書く

お名前:
URL:
コメント:
画像認証:
下の画像に表示されている文字を入力してください。

トラックバック 0

この広告は前回の更新から一定期間経過したブログに表示されています。更新すると自動で解除されます。