CUDA学习笔记(5) 原子操作

  原子操作是指，当一个线程（Thread）要对同一个显存变量依次进行“读-计算-写”的操作时，这个“读-计算-写”的操作必须连贯地执行，中间不能插入任何其他操作。

  举个例子，假设我们想要用GPU统计“char data_0[32] = {1,0, … ,1}”这个数组中“0”和“1”的个数并写入“int counter[2]”中。

  如果我们不使用原子操作，直接在核函数中这样写：

extern "C" __global__ void kernel_func(int * counter, char * data_0)
{
    // 计算线程号
    unsigned int block_index = blockIdx.x + blockIdx.y * gridDim.x + blockIdx.z * gridDim.x * gridDim.y;
    unsigned int thread_index = block_index * blockDim.x * blockDim.y * blockDim.z + \
        threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.x * blockDim.y;

    // 统计结果
    int value = data_0[thread_index];
    counter[value] ++;
}

  我们会发现结果是“counter[2] = {1, 1}”，这显然不是正确的结果。

  在GPU中，线程在并行时，每个线程（Thread）都会有一组专供自己的用的寄存器。例如上面的代码中，我们在核函数中定义了“block_index”、“thread_index”、“value”三个寄存器，此外编译器还会自动生成一个“counter”寄存器。也就是说，当这32个线程（Thread）并行执行的时候，会占用$4 \times 32$个32位寄存器。

*注1：自动生成一个“counter”寄存器的原因和CPU中汇编的的原理相同，由于“counter[value]”是存放在显存中的变量，对他进行“++”操作时会先将它读取到寄存器（Register）中，对这个寄存器变量进行操作后再将其写入原来的显存变量。如此便会在执行过程中占用一个寄存器。

  原因是这样的，在GPU中每32个线程（Thread）作为一个线程束（Warp）整体执行一系列操作。上图中的执行过程是这样的：

1. 线程束0（即线程0到31），从全局显存中读取了数组“data_0”的数值放在每个线程对应的寄存器“value”中。
2. 线程0到31几乎同时分别占用一个SP，然后第$i$条线程根据各自的寄存器“value”中的数值，准备读取显存变量“counter[value]”的值。（此时“counter[2] = {0, 0}”）
3. 线程束0读取显存变量“counter[value]”到各个线程的寄存器“counter”中。
4. 线程0到31同时对自己的“value”寄存器中的数值执行“++”。
5. 线程束0将寄存器“counter”值写入显存变量“counter[value]”中。（此时“counter[2] = {1, 1}”）

  因此，得到的结果是“counter[2] = {1, 1}”。这样，我们可以看到，错误实际上是出在了并行上——线程0还没有将自己计算的“counter”写回显存变量“counter[value]”，其他线程就已经读取了显存变量“counter[value]”的值。

*注2：上面的例子作为简单情况分析，例中只有32条线程（Thread），其数量小于空闲的流处理器（SP）数量（我的电脑上384个）时是这样的结果，每个线程都由一个流处理器（SP）来处理。在线程较多时可能多个线程都由一个流处理器（SP）处理。

  正确的方法是使用原子操作，在核函数中这样写：

extern "C" __global__ void kernel_func(int * counter, char * data_0)
{
    // 计算线程号
    unsigned int block_index = blockIdx.x + blockIdx.y * gridDim.x + blockIdx.z * gridDim.x * gridDim.y;
    unsigned int thread_index = block_index * blockDim.x * blockDim.y * blockDim.z + \
        threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.x * blockDim.y;

    // 统计结果
    int value = data_0[thread_index];
    atomicAdd(&counter[value], 1);
}

  注意“atomicAdd”函数，它的作用在于当线程$k$要操作显存变量“counter[0]”的值时，其他的线程若需要操作“c