OS XのOpenCL - その16

「プリンセス・プリキュア」のエンディングの曲とダンスが新しくなってた。リアルに近いシェーディングになってる。本編とのギャップがさらに広がった感じがする。それとも本編をもっとCGっぽい絵にするつもりなんだろうか。

それはいいとして、こないだの続き。パフォーマンスを測定する前に、まず基準を作る話。

13.4  計算/メモリアクセスのピーク性能ベンチマークを作る

コードの最適化を始める前に、特定のGPUデバイスでメモリをアクセスする場合と浮動小数点計算の場合にどのくらい速いかを見積もっておく必要がある。つぎの簡単なカーンルがベンチマークに使える。

コピーカーネル

つぎのリストのカーネルは大きな画像（このチャプタで示される16メガピクセルの画像）をインプットバッファから読んでアウトプットバッファにコピーする。GPUが可能な一番速いメモリアクセスのよい指標になる

kernel void copy(global const float * in,
                 global float * out,
                 int w,int h)
{
    int x = get_global_id(0); // ワークアイテム内の
    int y = get_global_id(1); // ピクセルの位置 = (x,y)
    out[x+y*w] = in[x+y*w]; // 読み込んで書き出す
}

図-10はイメージサイズに対するコピーカーネルのスピードを表している。

この図の場合、1秒間に13.5Gピクセルコピーできるということがわかる。これはそのデバイスの最速のメモリアクセスのよい指標である

MADカーネル

つぎのリストのような積和カーネル（Multply ADd、MAD）では、大きな画像を読み込んで画像データに対してある浮動小数点の計算をして、結果をアウトプットバッファに書き出す。MADカーネルはGPUの最大の計算スピードを見積もるのに使うことができる。しかしMADカーネルを使うときはメモリの制約が浮動小数点計算の能力をマスクしないということが保証できるようにたくさんの計算をする必要がある

kernel void mad3(global const float * in,
                 global float * out,
                 int w,int h) {
    int x = get_global_id(0);
    int y = get_global_id(1);
    float a = in[x+y*w];           // 読み込み
    float b = 3.9f * a * (1.0f-a); // みっつの浮動小数点演算
    out[x+y*w] = b;                // 書き出し
}

MADベンチマークは一連の依存性のある計算の場合の計算とメモリの比を示している。計算主体のカーネルは幾つかの独立した部分（依存性のある計算）含んでいるとき、かなり高い計算/メモリ比率に到達し得る。例えばマトリクスマトリクス積カーネルは同じGPUで2Tflops/秒に近い計算スピードが出る。図-11に示すように、3つの浮動小数点演算をコピーカーネルのコードに追加したとき（上にある赤い線）、

まだ11.9Gピクセル/秒の能力を持っている。これは3フロップス程度ではメモリ律速である、ということを示している。図-12はコピーカーネルに6つの浮動小数点演算を追加したとき（上の赤い線）、

まだ11.8Gピクセル/秒で処理できている。しなわち6フロップスでもまたメモリ律速である。 図-13はコピーカーネルに24演算を追加したとき、10.1Gピクセル/秒に低下していることを示す。

処理速度の低下するということは影響が十分大きいと言えるので、この結果はこのカーネルがGPUの計算のいいベンチマークとなっているということを表している