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

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

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

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

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

但是这个结构是存在问题的，我在测试时候发现调用theKernel(lock)可以正常运行，而theKernel(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中仅有一个线程。

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

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

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

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

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

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

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