《CUDA编程》9.原子函数的合理使用

在 CUDA 中，一个线程的原子操作可以在不受其他线程的任何操作的影响下完成对某个（全局内存或共享内存中的）数据的一套"读-改-写"操作。

0 基本概念

0.1 原子操作

原子操作是一种不可分割的操作，这意味着在执行该操作时，不会被其他线程中断。原子操作确保了对共享数据的更新是安全的，即使在多个线程同时尝试更新该数据时。

0.2 原子函数

原子函数是CUDA提供的一组内置函数，这些函数执行原子操作。

1 完全在GPU中进行归约

我们再次观察《CUDA编程》8.共享内存的合理使用中的归约代码，会发现并没有在GPU中做全部的计算，而是将数组返回到主机函数中进行最后结果的计算，为什么不直接在核函数中完成计算呢？原因如下：

• 线程块之间无法同步： 在核函数中，每个线程块只能处理自己分配到的数据，无法跨线程块进行全局的归约求和，所以只能放回主机函数中。

如果能够直接在GPU中完成最后的求和计算，只把结果传回主机，则或许可以更一步提升性能，目前有两种方法：

1. 使用另一个核函数对结果数组进行求和
2. 在原本的核函数末尾，利用原子函数进行归约求和，本章讨论该种方法

1.1 使用原子函数进行归约

观察原核函数末尾的代码：

if (tid == 0)
{
    d_re[blockIdx.x] = x[0];
}

即将每一个线程块中归约的结果从共享内存x[0]复制到全局内存 d_re[blockIdx.x]。如果想将不同线程块归约好的部分进行累加，理论上代码可以修改如下（但其实是无法按预期执行的）：

if (tid == 0)
{
	d_re[0] += x[0];
}

无法执行的是因为执行的次序无法确定，即数据竞争。因为线程块之间是并行运行的，如果多个线程块中的线程都试图同时修改d_re[0]，则会让结果变得无法预测，这时就应该使用原子函数：

// 使用原子加，将每个线程块的结果加到全局结果中
if (tid == 0)
{
    atomicAdd(&d_re[0], x[0]);
    //atomicAdd(address, val)
}

该函数的作用是将地址 address 中的旧值 old 读出，计算 old + val， 然后将计算的值存入地址 address。

这些操作在一次原子事务（atomic transaction）中完成， 不会被别的线程中的原子操作所干扰。原子函数不能保证各个线程的执行具有特定的次序， 但是能够保证不被其他线程干扰，所以能够保证得到正确的结果。

2 原子函数

原子函数对它的第一个参数指向的数据进行一次"读-改-写"的原子操作，即一气呵 成、不可分割的操作。第一个参数可以指向全局内存，也可以指向共享内存。

对所有参与的 线程来说，该"读-改-写"的原子操作是一个线程一个线程轮流做的，但没有明确的次序。 另外，原子函数没有同步功能。

2.1 所有的原子函数

下面，我们列出所有原子函数的原型，我们约定，对每一个线程 来说，address 所指变量的值在实施与该线程对应的原子函数前为 old，在实施与该线程 对应的原子函数后为 new。

对每一个原子函数来说，返回值都是 old。注意，原子函数只能在核函数中使用。

1. 加法：T atomicAdd(T *address, T val)
   功能：new = old + val
2. 减法：T atomicSub(T *address, T val)
   功能：new = old - val
3. 交换：T atomicExch(T *address, T val)
   功能：new = val
4. 最小值：T atomicMin(T *address, T val)
   功能：new = (old < val) ? old :val
5. 最大值：T atomicMax(T *address, T val)
   功能：new = (old > val) ? old :val
6. 自增：T atomicInc(T *address, T val)
7. 功能：new = (old >= val) ? 0 : (old+ 1)
8. 自减：T atomicDec(T *address, T val)
   功能：new = ((old == 0) || (old >val)) ? val : (old - 1)
9. 比较-交换（Compare And Swap)：T atomicCAS(T *address, T compare, T val)
   功能：new = (old == compare) ? val : old
10. 按位与：T atomicAnd(T *address, T val);
    功能：new = old & val
11. 按位或：T atomicOr(T *address, T val);
    功能：new = old | val
12. 按位异或：T atomicXor(T *address, T val)
    功能：new = old ^ val

T 表示相关变量的数据类型，各个原子函数对数据类型的支持情况见下表：