cuda-tutorial

5.1 Atomic演算

CUDAは、1つのスレッドが、あるアドレスに対して処理を行っている間に、他のスレッドに割り込まれない、排他的な処理をサポートしており、それをAtomic演算と呼ぶ。 Atomic演算を行う関数は複数ある[1]が、総和計算では atomicAdd() を用いるとよい。

Atomic演算を活用した総和計算は次の通り。

"sum_atomic.cu"

...

__global__ gpu_sum(int *sum_dev, int *arr_dev){
    int i_global = blockIdx.x * blockDim.x + threadIdx.x;

    if(i_global < Narr){
        atomicAdd(sum_dev, arr_dev[i_global]);
    }
}

...

sum_dev というアドレスに対して、arr_dev[i_global]という値を足し込んでいる。 こうすることでアクセスの衝突が起こらず、正しい動作をする。

しかし、これでは逐次処理を行っているのと大差なく、高速に動作するわけではない。 sum_dev にあるスレッドが書き込んでいる間、他のスレッドは待機を余儀なくされるからだ。 GPUの資源を活用するためにはもうひと工夫が必要だ。

Links

• 次へ
• 前へ
• 目次

引用、紹介

[1] CUDAアトミック関数