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中仅有一个线程。

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

方法2.通用原子操作

考虑到同一个wrap的线程都是‘同进退共生死’，那么我们只能在那个获得锁的线程退出函数前，就释放了锁。看代码：

__global__ void kernel1(){int index = 0;int mSize = 1;bool blocked = true;while(blocked) {if(0 == atomicCAS(&mLock, 0, 1)) {index = mSize++;doCriticJob();atomicExch(&mLock, 0);blocked = false;}}
}
int main(){kernel1<<<4,128>>>();cudaDeviceSynchronize();
}

在程序中，获得锁的线程进入到if中，并在执行完if之前就释放了锁，这样就解决了同一个wrap出现死锁的情况。当然，这样的写法不怎么优美且不鲁棒…（但是能用）。另外，这个函数换成这样写法就不行了：

__global__ void kernel2(){int index = 0;int mSize = 1;while(true) {if(0 == atomicCAS(&mLock, 0, 1)) {index = mSize++;doCriticJob();atomicExch(&mLock, 0);break}}
}
int main(){kernel2<<<4,128>>>();cudaDeviceSynchronize();
}

这是因为break在不同的机器和编译器中，不能都保证是先释放了锁再break出来,可能被编译器优化成其他形式。可以看出这种cuda通用原子操作确实比较蛋疼。

不过我在项目中采取了这种方法，将mLock由int变为int数组，就可以实现多把锁并行，提高效率，贴上我运行ok的代码：

__device__ void doCriticJob(int thread_index, float* mProcess) {mProcess[thread_index] += 0.222;printf("thread is: %d \n", threadIdx.x);
}
__global__ void kernel2(int* mFlag, float* mProcess) {bool blocked = true;int thread_index = (threadIdx.x + blockDim.x*blockIdx.x) % 4;while (blocked) {if (0 == atomicCAS(&mFlag[thread_index], 0, 1)) {doCriticJob(thread_index, mProcess);atomicExch(&mFlag[thread_index], 0);blocked = false;}}
}int main() {cudaError_t cudaStatus;float h_Process[4] = {0};int h_Flag[4] = {0};float *dev_Process;int *dev_Flag;cudaStatus = cudaMalloc((void **)&dev_Process, 4*sizeof(float));if(cudaStatus != cudaSuccess){ fprintf(stderr,"malloc 1 failed\n");}cudaStatus = cudaMalloc((void **)&dev_Flag, 4*sizeof(int));if(cudaStatus != cudaSuccess){ fprintf(stderr,"malloc 2 failed\n");}cudaStatus = cudaMemcpy(dev_Process, h_Process, 4*sizeof(float), cudaMemcpyHostToDevice);if(cudaStatus != cudaSuccess){ fprintf(stderr,"malloc 3 failed\n");}cudaStatus = cudaMemcpy(dev_Flag, h_Flag, 4*sizeof(int), cudaMemcpyHostToDevice);if(cudaStatus != cudaSuccess){ fprintf(stderr,"malloc 4 failed\n");}kernel2<<<2, 3>>>(dev_Flag,dev_Process);cudaStatus = cudaDeviceSynchronize();if (cudaStatus != cudaSuccess) {fprintf(stderr, "ffffff is %d\n",cudaStatus);}float outProcess[4];cudaStatus = cudaMemcpy(outProcess, dev_Process, 4*sizeof(float), cudaMemcpyDeviceToHost);if (cudaStatus != cudaSuccess) {fprintf(stderr, "kkkkkk is %d\n",cudaStatus);}for (float mProces : outProcess) {std::cout << mProces << std::endl;}
}

可以实现对多个锁的控制，锁之间是并行的。

方法3.仲裁中介

这个方法的名字是我瞎取的。直接上代码：

__device__ volatile int sem = 0;__device__ void acquire_semaphore(volatile int *lock){while (atomicCAS((int *)lock, 0, 1) != 0);}__device__ void release_semaphore(volatile int *lock){*lock = 0;__threadfence();}..........__global__ void inKernel(){...__syncthreads();if (threadIdx.x == 0)acquire_semaphore(&sem);__syncthreads();//begin critical section// ... your critical section code goes here//end critical section__threadfence(); // not strictly necessary for the lock, //but to make any global updates in the critical //section visible to other threads in the grid__syncthreads();if (threadIdx.x == 0)release_semaphore(&sem);__syncthreads();...
}

为什么叫仲裁中介呢？因为这里采用了每个block的一个线程作为中介进行仲裁，在acquire_semaphore中争夺锁，一旦某个block的第一个线程获得了锁，那么剩下的block第一个线程将陷入while循环中，同时因为__syncthreads()，导致整个block停下。这就实现了以block为单位的串行。但是这其实也不完美，若要再进一步在block的线程中实现串行，则要继续加入条件判断。可以看StackOverflow：链接

总结：

cuda 中不可避免的遇到需要串行计算的情况，可以每个方案都不是完美的，需要根据情况进行取舍，我也继续学习，应该是存在更好的方案，日后遇到了再添加进来。