SSブログ

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

今朝はさすがに地下鉄がすいてた。みんな休んでるんだろうなあ。今は設計部部長さんの作業待ち。それが終わると僕の光学測定に入る。

ということでその間に、抄訳の続き。前回ガウシアンブラーのナイーブな実装に対して、効率を上げるためにメモリアクセスの工夫を導入した。今日のところはその改善の度合いと、さらなる高速化のための追加の作業。最終的には前回の最初の状態にくらべて一桁以上速くなる、というもの。

ところが残念ながらこの(前回転置を行った上で、処理対象となるメモリのブロックをローカルメモリに移動させる)コードはあまり速くならない。
0810fig28.png
問題はローカルメモリの横の列からグローバルメモリの縦の列にコピーするときにこれまでと違うメモリアクセスのパターンを作ってしまったことである。
0810fig29.png
これを解決するためには、対角方向に画素をコピーするようにワークグループをマップするように変更する。
0810fig30.png
コードを入出力のためのコピーを斜めに実行するように変更するには、1行書き換えればいい。
kernel void transposeLS(global const float * in,
                        global float * out,
                        int w,int h)
{
    local float aux[256];           // ブロックサイズは16x16
    int bx = get_group_id(0),       // (bx,by) = 入力ブロック
    by = get_group_id(1);
    int ix = get_local_id(0),       // (ix,iy) = ブロック内の一つの画素
    iy = get_local_id(1);
    // 変更した行
    by = (by+bx)%get_num_groups(1); // 斜めマッピング
    
    in += (bx*16)+(by*16)*w; // 入出力の原点に移動
    out += (by*16)+(bx*16)*h;
    aux[iy+ix*16] = in[ix+w*iy];    // ブロックの読み込み
    barrier(CLK_LOCAL_MEM_FENCE);   // 同期
    out[ix+h*iy] = aux[ix+iy*16];   // ブロックの書き出し
}
このバージョンは速くなる。
0810fig31.png
このローカルメモリの転置のコードを使うとガウシアンブラーはすごく速くなる。
0810fig32.png
まだ、処理は最初の見積ほどは速くはない。問題は再帰ガウシアンのループはもともとシーケンシャルなものなので、GPUのワークグループを十分に飽和させることができないという点である。最初の見積までパフォーマンスをあげるためには並列化のレベルを上げるようなアルゴリズムに変更する必要がある。

13.8  GPUでのオーバーヘッドを減らすためのガイドライン

GPUで実行するOpenCLの効率を改善するための一般的な原理。
  • OpenCLのプログラムをビルドするのは計算資源的に高価で、プロセスに一回だけにすべきである。
    一度だけコンパイルして何回も実行するOS X v10.7以降のツールが利用できる。実行中にカーネルをコンパイルすることを選択したなら、そのコストに見合うだけそのカーネルを何度も実行しなければいけない。最初にプログラムが実行されたあと、コンパイル済みのコードを次の実行時に再利用するためにバイナリを保存することができるが、OpenCLの改版やホストマシンのハードウェアの変更のせいでビルドが失敗するかもしれないのでカーネルの際コンパイルの準備をしておかないといけない。 また、ソースコードの代わりにOpenCLコンパイラが生成したビットコードを使うこともできる。コンパイル済みのビットコードを使えば処理速度を上げられるし、アプリと一緒にソースを出荷する必要性も減らすことができる
  • 可能な限りOpenCLのビルトイン関数を使うように。それらの関数を使えば最適化されたコードが生成できる
  • 精度の速度のバランスをとるように。GPUはグラフィクス用に設計されていて、精度に対する要求は低い。fast_、half_、native_などのOpenCLのビルトインのようにいくつかの変形がある。プログラムビルドオプションでいくつかの速度に対する最適化の制御ができる
  • 分岐を避けるように。GPUではすべてのスレッドが同じコードを同時に実行するようにスケジュールされる。その結果、条件文を実行するとき、すべてのスレッドは両方の分岐を実行し、不成立の分岐に対しては結果を書き出さない。条件文を使わない(a?x:y演算子に置き換える)、あるいはビルトイン関数を使うのが最も良い
  • バッファの代わりにイメージオブジェクトを使うように。ある場合において(特有のメモリアクセスパターンでは)イメージをアクセスするときにGPUが使うハードウェアデータパスは、バッファを使った場合より速いことがある。とくに16ビット浮動小数点データ(half)を使うときはバッファよりもイメージを使う方が重要である

14  CPUでのパフォーマンスの改善

{{decafish : ここはもういいや}}
nice!(0)  コメント(0)  トラックバック(0) 

nice! 0

コメント 0

コメントを書く

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

トラックバック 0

OS XのOpenCL - その19Windows10 on Raspber.. ブログトップ

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