CUDA原子操作

发布时间 2024-01-01 22:50:12作者: 牛犁heart

这节主要涉及到一个多线程情况下存在的数据竞争问题 -- 多个线程同时访问共享数据时,由于没有正确的同步机制,导致数据出现不一致的情况。

  • C/C++ 多线程中,可以通过互斥锁(mutex)、原子操作(atomic,C++11 也提供了原子操作库,如std::atomic,用于实现原子加法、原子赋值等操作)、线程局部变量(C++11提供了thread_local,它为每个线程分配了一份独立的内存,使得不同线程之间的局部变量互不干扰。这样可以避免多个线程访问同一共享数据时的数据竞争问题)以及使用同步代码块(using)、全局变量等措施来避免多线程操作

  • 由于CUDA采用单指令多线程(Single-Instruction Multiple-Thread,SIMT)来管理和执行 GPU 上的众多线程,因此在操作数据时,也会存在数据竞争等问题,而在CUDA中避免数据竞争的方法主要包括使用原子操作适当的索引计算

原子操作可以确保在多线程环境下对共享变量进行原子的读写操作,从而避免数据竞争。
利用线程索引计算可以使得每个线程操作不同的数据元素,避免竞争

上代码:

__global__ void increment_naive(int *g)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    idx = idx % ARRAYSIZE;
    // g[idx] = g[idx] + 1;
    printf("before %d:%d\n", idx, g[idx]);
    atomicAdd(& g[idx], 1);
    printf("after %d:%d\n", idx, g[idx]);
}

从主机代码中将数据拷贝到设备端,默认都是放在设备的全局内存中(也就是global memory中),在运行核函数时从全局内存中读数据、运算、写入全局内存

  • 首先多个线程会同时从全局内存中取出g[idx]的数据,
  • 运算, g[idx] = g[idx] + 1; 虽然此处计算了,也写入了对应的g[idx]中,但是其他线程并没有拿到这个线程计算的的结果,他们用的数据还是最开始的原始数据,导致最后计算出错
    atomicAdd(& g[idx], 1);该函数为CUDA提供的原子函数接口,该接口详细描述如下:
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);
__half2 atomicAdd(__half2 *address, __half2 val);
__half atomicAdd(__half *address, __half val);
__nv_bfloat162 atomicAdd(__nv_bfloat162 *address, __nv_bfloat162 val);
__nv_bfloat16 atomicAdd(__nv_bfloat16 *address, __nv_bfloat16 val);

上面是 device-wide atomicAdd 不同数据类型的定义,表示从第一个参数 address 指针指向的内存地址(可以是全局内存或共享内存)中读取16位、32位或64位数据(记做旧值 old),与第二个参数val 做一个加法计算(old + val),然后将求和结果写回到 address 指针指向的内存地址中,并返回未做计算前的旧值 old。这三个操作在一个原子事务中执行。

  • 32位 floating-point 浮点版本的 atomicAdd() 仅由计算能力 2.x 及更高版本的设备支持。
  • 64位 floating-point 浮点版本的 atomicAdd()只被具有计算能力 6.x 及更高版本的设备支持。
  • 32位 __half2 浮点版本的 atomicAdd() 只被具有计算能力 6.x 及更高版本的设备支持。对于两个__half 或__nv_bfloat16 元素,分别保证 __half2 或 __nv_bfloat162 添加操作的原子性;整个__half2 或 __nv_bfloat162 不能保证作为单个32位访问是原子的。
  • 16位 __half 浮点版本的 atomicAdd() 仅由计算能力为 7.x 及更高版本的设备支持。
  • 16位 __nv_bfloat16 浮点版本的 atomicAdd() 仅由计算能力为 8.x 及更高版本的设备支持。