rigayaの日記兼メモ帳

最後に、配列の総和計算でshared local memory(SLM)の使用方法を確認する。

---

配列の総和の計算は、CPUだと非常に単純で、

float sum = 0.0f;
for (int i = 0; i < nSize; i++) {
    sum += ptrArray[i];
}

とまあ、これだけなのだが、このままのループの形だと並列性がないので、GPUでやろうとするとやけに面倒だったりする。

こういったいわゆるreductionの計算は、なんとか並列性をひねりだし、work group内で共有できるメモリ(shared local memory, SLM)を使うと効率的に計算できることが知られている。

計算の仕方としては、もう図にしたほうが早いので、下のような感じ。(図はwork groupあたり16 work itemの例)
CUDAでもお馴染みの方法だったりする(最近はshuffleとかも使うらしいが)

まあ、ひとつづつ足すのではなく、なるべく並列に計算できるところを見つけて、GPUを活かそう、というもの。無限に演算器があるなら、計算のオーダーをO(n)からO(log(n))に落とすイメージ。

通常、GPUでは異なるスレッドの計算結果を得ることはできないが、__localメモリを使えば、work group内の計算結果を取得できることを利用している。



このようにkernelを1回起動することで、work group内のwork item分の和を取ることができる。つまり、work groupあたり16work itemの設定なら、1回で要素数を1/16にすることができ、これを繰り返すことで総和を得ることができる。もちろん、work groupあたりのwork item数が大きいほど一気に要素数を減らせるので、実際にはwork groupあたり16work itemなどにする。

実際のkernelコードはこんな感じ。

__kernel void reduce_add(__global float* pIn, __global float* pOut, int nSize) {
    const int lid = get_local_id(0);
    const int gid = get_global_id(0);
    __local float shared[GROUP_SIZE];

    shared[lid] = (gid < nSize) ? pIn[gid] : 0;
    barrier(CLK_LOCAL_MEM_FENCE);
    for (int offset = get_local_size(0) >> 1;
        offset > 0; offset >>= 1) {
        if (lid < offset) {
            shared[lid] += shared[lid + offset];
        }
        barrier(CLK_LOCAL_MEM_FENCE);
    }
    if (lid == 0) {
        pOut[get_group_id(0)] = shared[0];
    }
}

shared local memoryは、kernelコードでは__localで取ることができる。

__localに書き込んだ値を他のスレッドで使う場合、かならず

barrier(CLK_LOCAL_MEM_FENCE);

によりwork group内のスレッドの同期をとる必要がある。work group内のスレッドはすべてが同時に実行されるわけではないので、きちんとすべてのスレッドが書き込みを終えてから次に進むようにしないと、おかしな値を使用してしまう可能性があるからである。

実際に実行している様子はこちら。



よく見えないけど、reduction_addが2回実行されていて4M要素の総和を取っている。(4M → 16K → 64)

---

この方法、GPUでもある程度高速に総和を取れてよいのだけど、加算の計算順序が変わるので、計算結果がCPUと多少変わるのが面倒なところ。浮動小数点演算は数学の計算と違って可換ではないのだ…。

コードはこちら。

---

OpenCLは勉強自体はこれまで少しづつやっていたことだが、時間がなかったのでこれまでまとめたりできなかった。この3連休でやっとまとめることができたので、これできっともう忘れないだろう(棒)

だいたい通常の計算と、image、__localの使い方がわかったし、これでやっとOpenCLの基本の基本ぐらいはわかったように思う。おおまかにCUDAとの対応関係もつかめたし、多少はCUDAの知識も使えそうな感じ。

あとは、実際にやりたいこと次第、になるのだと思う。

---

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)の使用 (いまここ)

---