rigayaの日記兼メモ帳

今回は、この前の簡単な配列加算のプログラムの速度を確認し、work sizeの調整で実行速度が変わることを確認してみる。

---

OpenCLのkernelがどのように実行されているのか確認したい。

まあ、vtuneがあれば悩まなくていいのだが、さすがにあれはお値段がお高すぎる。とりあえずいろいろ調べると、簡単にはIntel Graphics Performance Analyzersを使えばよいようだ。

(なんか昔も調べた気がするので、たぶん調べるの2回目…)

まず、Intel Graphics Monitorを起動する。「Analyze Application」から、起動する実行ファイルのパスを指定する。



正しくプロファイルを取得するには、トレースする時間を適切に指定する必要がある。[Managae Profiles...]から、[Tracing]タブのTracing Durationを調整する。



以下のような順序になるよう調整しないと、うまくトレースが取得できなかった。

1. 起動
2. OpenCLが実行され、完了
3. 起動してから、Tracing Durationの時間が経過し、トレースが終了
4. プログラムが終了

特にはまるのは、3と4が逆転してしまうと、トレースが取得できなかったことで、プログラムのほうにOpenCL実行後、

Sleep(5000);

を入れて調整した。なんかもっといいやり方がありそうだけど…。

さて、取得したトレースはIntel Platform Analyzerで見ることができて、こんな感じ。



キューにタスクが積まれ、転送→転送→Add→転送の順序で実行している様子とその時間が確認できる。

---

さて、この前のkernelは、以下のような単純な配列加算を行うものだった。

__kernel void Add(__global int* pA, __global int* pB, __global int* pC) {
    const int x     = get_global_id(0); //0次元目についてglobal idを取得
    const int y     = get_global_id(1); //1次元目についてglobal idを取得
    const int width = get_global_size(0);

    const int id = y * width + x;

    pC[id] = pA[id] + pB[id];
}

このkernelを以下のように実行していた。

//work sizeの指定
//ここでは1要素に対して1 work item
//またグループあたり1 work item (実は効率的でない)
//width * heightの2次元でwork itemを作成
size_t globalWorkSize[2] = { width, height };
size_t localWorkSize[2] = { 1, 1 };

//タスクをキューに積む
clEnqueueNDRangeKernel(
    ocl->commandQueue, //タスクを投入するキュー
    ocl->kernel,       //実行するカーネル
    2,                 //work sizeの次元
    NULL,              //NULLを指定すること
    globalWorkSize,    //全スレッド数
    localWorkSize,     //1グループのスレッド数
    0,                 //この関数が待機すべきeventの数
    NULL,              //この関数が待機すべき関数のリストへのポインタ
    NULL);             //この関数の返すevent

このwork size指定では、1つのwork itemを持つwork groupがwidth * heightだけ作られることになる。しかし、前回見たGPUの構造から、この指定は以下の理由で効率が悪いことがわかる。

・複数のwork itemをまとめて実行するSIMD-32等が効果を発揮できない
同じwork groupのwork itemしかまとめることはできない

・異なるwork groupだとどのsubsliceで実行されるわからない
メモリアクセスがばらばらで非効率

そこで、通常第0次元のwork itemを複数まとめて実行するよう指定する。このとき、一般的に32, 64など2の累乗がよいとされる。また、OpenCL 1.xでは、globalWorkSizeはlocalWorkSizeで割り切れる必要があることに注意して、localWorkSizeの倍数に切り上げる。

ここでは、work groupあたり32のwork itemを実行するよう指定する。

int ceil_int_div(int i, int div) {
    return (i + div - 1) / div;
}

//iをdivで切り上げ
int ceil_int(int i, int div) {
    return ceil_int_div(i, div) * div;
}

cl_uint ExecuteAddKernel(ocl_args_d_t *ocl, cl_uint width, cl_uint height) {
    cl_int err = CL_SUCCESS;

    size_t localWorkSize[2]  = { 32, 1 };
    size_t globalWorkSize[2] = {
        ceil_int(width,  localWorkSize[0]),
        ceil_int(height, localWorkSize[1])
    };

    (...略...)

切り上げを行うことで、work item数が実際の数を上回ることがあり得る。確保していないメモリ領域にアクセスしないよう、kernelコードを以下のように変更して、widthとheightは引数で与えるようにする。

__kernel void Add(__global int* pA, __global int* pB, __global int* pC, int width, int height) {
    const int x      = get_global_id(0);
    const int y      = get_global_id(1);

    if (x < width && y < height) {
        const int id = y * width + x;
        pC[id] = pA[id] + pB[id];
    }
}

---

kenrnel実行の所要時間を確認する。

work groupあたり1つのwork item = 33.11ms




work groupあたり32のwork item = 1.64 ms



というわけで劇的に高速化していることが確認できる。このように、work sizeの調整は重要であるが、これをOpenCL runtimeに自動的に調整させることも可能である。その際は、clEnqueueNDRangeKernel()のlocal_work_sizeにNULLを指定すればよい。

---

今回の変更はこちら。

---

続き>>

---

OpenCLメモリスト
OpenCLメモ (その1) - かんたんな計算
OpenCLメモ (その2) - Intel GPUの構造
OpenCLメモ (その3) - work sizeの調整 (いまここ)
OpenCLメモ (その4) - 転送(コピー)の排除 (USE_HOST_PTR)
OpenCLメモ (その5) - 転送(コピー)の排除 (SVM : OpenCL 2.0)
OpenCLメモ (その6) - imageの使用
OpenCLメモ (その7) [終] - reductionとshared local memory(SLM)の使用

---