cuda 原子锁&多线程操作&通用原子操作

经历了很多奇奇怪怪的bug，整理一下。先描述要做的事情以及怎么做：

在项目中，空间中有200w+的点，需要映射到一个grid_map的600*600的网格中，落入到同一个格子的点需要进行一些计算获得一个值。对于格子与格子之间是并行的，但格子之中的点需要设计为串行。所以在计算某个格子中的点时，需要将格子的值保护起来，只允许一个线程（点）计算并改变。

这里就用到了cuda的通用原子操作。也许有人会问，cuda提供了一些原子操作函数，能不能直接用呢？cuda提供的原子函数适用于简单的单一变量判断加减，而对于需要复杂的计算操作是力不从心的。但其实，我们要实现的通用原子操作也是基于cuda的原子函数，我们进行一些设计就可以得到想要的通用原子操作，比如锁。

方法1.原子锁

在《GPU高性能编程CUDA实战》一书中，提到了通用原子操作的锁的设计，贴上源码：

struct Lock {
   
    int *mutex;
    Lock(void) {
   
        int state = 0;
        cudaMalloc((void **) &mutex, sizeof(int));
        cudaMemcpy(mutex, &state, sizeof(int), cudaMemcpyHostToDevice);
    }
    ~Lock(void) {
   
        cudaFree(mutex);
    }
    __device__ void lock(void) {
   
        while (atomicCAS(mutex, 0, 1) != 0);
    }
    __device__ void unlock(void) {
   
        atomicExch(mutex, 0);
    }
};
.......

__global__ void theKernel(Lock myLock) {
   
	myLock.lock();
	Do_your_job();
	myLock.unlock();
}

这里通过atomicCAS和atomicExch两个函数进行设计，但一个线程lock之后，将mutex置为1，其他线程将在while处循环等待，直到该线程unlock，将mutex重新置于0，剩下的线程中再次争夺锁。

但是这个结构是存在问题的，我在测试时候发现调用theKernel<<<128, 1>>>(lock)可以正常运行，而theKernel<<<1, 128>>>(lock)出现了死锁，也就是在block中线程数大于1情况中，出现死锁。百思不得其解…后来查到了出现这种情况的原因：

cuda运行是以wrap为单位进行的，也就是说一个wrap中32个线程中的一个获得了锁，执行完了lock，按理说该线程要继续执行Do_your_job()再unlock，而现实是线程都卡在了lock处。这就是因为wrap的同步执行规则(locked-step execution)，换句话说，一个wrap的线程是同步执行一个函数，并同步退出一个函数。获得锁的线程在lock函数结束处苦苦等待其他31个线程兄弟一起进入Do_your_job()，而剩下的31个线程却等着它unlock释放锁，所以出现了死锁。而每个block中只有一个线程则不会出现死锁，是因为此时wrap中仅有一个线程。

显然，这个设计方法并不满足我的需求。