最近在研究CUDA。其中有一个需求是要对一段代码加锁。CUDA里只提供了现成的对单个变量进行操作的原子操作库[1]，而没有提供对一段代码加锁的机制，于是我就想怎么改造一下这个原子库。
一开始自己写了一段代码，差不多是这个样子的[2]。
1 2 3 4 5 6 | __device__ inline void lock(int* p_mutex ){ while(atomicCAS(p_mutex, 0, 1) != 0); // compare and swap } __device__ inline void unlock(int* p_mutex ){ atomicExch(p_mutex, 0);// exchange } |
调用了之后发现从加锁的循环里压根就跑不出来。自己想了半天死活也搞不懂。然后从StackOverFlow上查到资料。发现这个和CUDA的Warp有关[3]。每个Warp由32个线程组成，它们必须同时执行相同的指令。如果用上述代码实现自旋锁。假定Warp中有一个Thread获得了锁，需要向下执行才能够释放锁。但是其他Thread却没得到锁，因此整个Warp就不能向下执行，始终卡在while循环里形成死锁。这种死锁叫SIMD死锁(SIMD DeadLock)。
之后翻了一大堆文献和网站，里面成堆的人都是实现同上，没有一个能用的。最后翻到了这篇文章[4]，解决了问题。整理出的自旋锁的模板如下：