cuda 通用原子操作

在多线程中，访问共享资源时，确保所有其他的线程都不在同一时间内访问相同的资源，就叫作原子性。在主机上，一般通过互斥锁（mutex,lock)的方式来保证线程之间的原子性。那么在cuda中，如何保证核并行的原子性？

cuda对一些常用操作包括加、减、按位与、或等提供了原子操作的函数，具体可参考博客https://blog.youkuaiyun.com/dcrmg/article/details/54959306。但cuda没有提供比较通用的互斥、锁方式，对于其他操作，要保证原子性，只能自己模拟这种互斥锁行为。

互斥锁的基本思想并不复杂，如我们可以拿一个标志位flag,如初始值赋为0，当有线程跑到flag时，看它为0，说明没有线程使用下面的操作，然后把flag置为1，表示此路不同，当要锁的操作完成后，flag=0,即解锁，示例代码如下：

__global__ kernel()
{
   ......
   if(flag==0)
   {
      flag=1; //加锁
      ...... //原子操作
      flag=0; //释放锁
      ......  
   }
}

当然这段代码是有问题的，因为if语句处并没有锁，多个线程可以同时读到if(flag=0)之后，flag=1之前，多个线程实质已经进入if语句内，这仍然不可避免造成非原子操作。因此简单的标志位赋值是不行的，要保证标志位的判断修改也要执行原子操作，幸运的是，cuda提供了原子类函数atomicCAS()可以实现上述形式，首先看非原子CAS操作如下：

__device__ int CAS(int flag,int compare,int val)
{ 
   int old=flag;
   flag=compare?val:compare;
   return old;
}

__global__ kernel()
{
   ......
   while(CAS(flag,0,1)!=0);
   .......//原子操作
   flag=0;
}

将上述程序中CAS改为atomicCAS的形式，就能实现GPU核内互斥锁的功能。

为了增加可读性，在《GPU高性能编程CUDA实战》一书中定义了锁的结构体：

struct Lock
{
  int *mutex;
  Lock()
  {
     int state=0;
     cudaMalloc((void**)&mutex,sizeof(int));
     cudaMemcpy(mutex,&state,sizeof(int),cudaMemcpyHostToDevice);
  }
 ~Lock()
  {
     cudaFree(mutex);
  }
 __device__ void lock()
  {
     while(atomicCAS(mutex,0,1)!=0);
  }
 __device__ void unlock()
  {
     atomicExch(mutex,0);//*mutex=0的原子操作，这里是为了增加可读性，直接*mutex=0也不会造成不安全操作
  }

__global__ void kernel(Lock lock)
  {
    ......
    lock.lock();
    ......
    lock.unlock();
    ......
  }