分类目录归档:Cuda

CUDA上的自旋锁(SpinLock)实现问题

最近在研究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],解决了问题。整理出的自旋锁的模板如下:

继续阅读