第二回です。今回はGPUでRadix sortをやる上で重要な処理である、累積和についてだけ説明します。
累積和とは
概要は前回時期を参照してもらいたいのですが、配列要素を$a_n$とした時の累積和を$S_n$は
$$S_n = a_0 + a_1 +\cdots + a_{n-1}$$
とします。$a_n$を和に含める定義もありますが今回はRadix sortをするので上の定義でいきます。
GPUの仕組み
GPUの並列処理ではgrid, block, threadの概念を理解しなければいけません。簡単に言うとthreadは並列化の単位で、各スレッドで処理が回ります。そしてblockはスレッドの集合体で、gridはblockの集合体です。カーネルを実行する際は、配列のサイズによってgridとblockをいくつ使うかを指定するわけです。こんな感じで
int block = 512;
// int の割り算ではあまりを切り捨てるので、配列サイズによってはgrid = 0になってしまいます。
// それを防ぐために以下のような式にしています。
int grid = (/*array_size*/ + block - 1) / block;
myKernel<<<grid, block>>>(/*variables*/);
ややこしいのは、gridはいくつのblockを使うか、blockは各blockでいくつのthreadを使うかを指定しているところです。gridとblockと書いていますが、意味的にはblockとthreadの数ですね。なんでこんなややこしいんだ。。。
GPUにカーネル実行の命令を投げたとき、設定したgrid, blockに応じてスレッドを割り振ります。このとき、各スレッドは完全に並行処理されるわけではなく、自動でスケジューリングされて時間差で実行されます。
ですので、他のスレッドの処理結果を自分のスレッドで使おうとすると、相手のスレッドがまだ処理完了していないと言う事態が発生します。
ここで同期が重要になってきます。他のスレッドの結果を利用する前に、全員の処理が終わるまで待機します。CUDAの場合、同期の種類は主に3つあります。
| 種類 | 関数 | 呼び出す場所 |
|---|---|---|
| block内だけの同期 | __syncthreads() | カーネル関数内 |
| block間の同期 | 特にない。カーネルを抜けることで同期する。 | - |
| CPUとGPU間の同期 | cudaDeviceSyncronize() | CPU側のホストコード内 |
素人考えだとblock間の同期とCPUとの同期があれば十分に思えますが、block間の同期は全blockを待機するのに対し、syncthreads()はblock内でだけ待機すればいいので、待機時間が短く高速です。cudaの高速化にはsyncthreads()をいかに上手く使うかが肝になります。
GPUで累積和をとる。
やっと累積和の話に移ります。配列数がNの時、cpuなら$O(N)$の計算量です。ところがどっこいGPUならこんな形で$O(\log(N))$にすることができます。(GPUの方がクロックが遅いことには留意しましょう。)
しかしこれはあまり効率的ではありません。たとえば$x_0$から$x_3$をblock0で、$x_4$から$x_7$をblock1で、実行するとき、d=2の時にいきなりblock間の同期が必要になります。$\log(8) = 3$ステップで処理ができるのはいいのですが、block間同期のせいで遅くなります。できるだけblock内の同期で済ませるため、和をとる処理を次のように分けます。
upsweep
downsweep
よくもまあこんなフローを思い付いたもんだと感心します。upsweepを見ると、d=2を実行する時まではblockを跨いだ配列のやりとりがないので、block内の同期で事足ります。downsweepもblock間の同期が削減できます。
つぶやき
2日連続で雑多な記事を書いたので追記することがあればまた書こうと思います。とりあえずRadix sortのGPU実装は以上です。