GPU编程自学9 —— 原子操作

深度学习的兴起，使得多线程以及GPU编程逐渐成为算法工程师无法规避的问题。这里主要记录自己的GPU自学历程。

目录

• 《GPU编程自学1 —— 引言》
• 《GPU编程自学2 —— CUDA环境配置》
• 《GPU编程自学3 —— CUDA程序初探》
• 《GPU编程自学4 —— CUDA核函数运行参数》
• 《GPU编程自学5 —— 线程协作》
• 《GPU编程自学6 —— 函数与变量类型限定符》
• 《GPU编程自学7 —— 常量内存与事件》
• 《GPU编程自学8 —— 纹理内存》
• 《GPU编程自学9 —— 原子操作》
• 《GPU编程自学10 —— 流并行》

九、 原子操作

原子操作 是指对全局和共享内存中的32位或者64位数据进行 “读取-修改-覆写”这一操作。

原子操作可以看作是一种最小单位的执行过程。 在其执行过程中，不允许其他并行线程对该变量进行读取和写入的操作。 如果发生竞争，则其他线程必须等待。

下面先给出原子操作函数的列表，后续会给出一个应用例子。

9.1 原子操作函数列表

9.1.1 atomicAdd()

int atomicAdd(int* address, int val);
unsigned int atomicAdd(unsigned int* address, unsigned int val);
unsigned long long int atomicAdd(unsigned long long int* address, unsigned long long int val);
float atomicAdd(float* address, float val);
double atomicAdd(double* address, double val);

读取位于全局或共享存储器中地址address处的32位或64位字old，计算(old + val)，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。

注意：

• 32位浮点数的操作只适用于计算能力大于2.0的GPU
• 64位浮点数的操作只适用于计算能力大于6.0的GPU

但可以通过以下操作在计算能力不足的GPU上实现浮点数原子操作：

#if __CUDA_ARCH__ < 600 
__device__ double atomicAdd(double* address, double val) 
{ 
    unsigned long long int* address_as_ull = (unsigned long long int*)address; unsigned long long int old = *address_as_ull, assumed; 
    do { 
        assumed = old; 
        old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val + __longlong_as_double(assumed))); 
        // Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN) 
        } 
        while (assumed != old); 
        return __longlong_as_double(old); 
} 
#endif