CUDAの累積集計方法

1974 ワード

1つのn要素配列の中の値をすべて加算するプログラムは、みんなが書いて、1つのforループで終わります.しかし、cudaコードにもそう書けば、cpuよりも速度が遅くなります.結局、人が一対一でつづると、cpuはできないに違いありません.
cudaにおけるアルゴリズムはn個のスレッドが共にこの演算に関与する.
nvidiaには公式のアルゴリズムがあるようで、ほとんどのcudaコードで使われています.

for(int i=(blockDim.x>>1);i>0;i>>=1){
    if(threadIdx.x < i){
        pis[threadIdx.x] += pis[threadIdx.x+i];
    }
    __syncthreads();
}

if(threadIdx.x == 0)    r[0] = pis[0];

配列の内容をsharedメモリに対応付けて読み,二叉を並べ替える.
しかし,このアルゴリズムはnが2のべき乗の場合のみをサポートし,そうでなければ多くの要素が漏れる.例えばn=3で、3番目の要素が漏れてしまいます.
しかし、単一スレッドアルゴリズムに復元するのは、本当に遅くて、とても悔しいです.新しいアルゴリズムを考えた.

for(int i = 1;i < blockDim.x ; i <<= 1){
    if(threadIdx.x % (i<<1) == i){
        tmp[threadIdx.x - i] += tmp[threadIdx.x];
    }
    __syncthreads();        
}    

if(threadIdx.x == 0)    r[0] = tmp[0];

nが2のべき乗でない場合も通常通り動作し,累積が必要なのはスレッドの一部である場合,アルゴリズムの前に超過したスレッドreturnを落とすだけでよい.ロバスト性がとてもいいです.原理は簡単ですが、抽象的で、基本的に奇数を前の位置に積み重ねて、結果の偶数たちとして新しい配列を形成して、続けます.0しか残っていないまで.
アルゴリズムの複雑さは両方ともlgNであり,(Nがspの数より少ない場合),新しいアルゴリズムはモデリングアルゴリズムに時間がかかるはずであるため,少し遅いはずである.オリジナルですが、前にも書いた人がいるに違いありませんが、しばらくはgoogleにはなりません.Googleが見つけられるように記録してください.

バンプ-graham-scanアルゴリズム

bit 1088