对于我正在编写的教程,我正在寻找一个“现实”且简单的死锁示例,该示例是由于对 SIMT/SIMD 的无知而导致的。
我想出了这个片段,这似乎是一个很好的例子。
我们将不胜感激。
…
int x = threadID / 2;
if (threadID > x) {
value[threadID] = 42;
barrier();
}
else {
value2[threadID/2] = 13
barrier();
}
result = value[threadID/2] + value2[threadID/2];
我知道,它既不是正确的 CUDA C 也不是 OpenCL C。
最佳答案
一个简单的死锁实际上很容易被新手 CUDA 程序员捕捉到,当一个人试图为单个线程实现一个临界区时,它最终应该由所有线程执行。它或多或少是这样的:
__global__ kernel() {
__shared__ int semaphore;
semaphore=0;
__syncthreads();
while (true) {
int prev=atomicCAS(&semaphore,0,1);
if (prev==0) {
//critical section
semaphore=0;
break;
}
}
}
atomicCAS
指令确保 exaclty 一个线程获得 0 分配给 prev,而所有其他线程获得 1。当一个线程完成其临界区时,它将信号量设置回 0,以便其他线程有机会进入临界区。
问题是,当 1 个线程获得 prev=0 时,属于同一个 SIMD 单元的 31 个线程获得值 1。在 if 语句中,CUDA 调度程序将单个线程置于暂停状态(将其屏蔽)并且让其他 31 线程继续工作。在正常情况下,这是一个很好的策略,但在这种特殊情况下,您最终会得到 1 个永远不会执行的临界区线程和 31 个等待无穷大的线程。死锁。
还要注意,break
的存在将控制流引导到 while
循环之外。如果你省略了 break 指令,并且在 if-block 之后还有一些应该由所有线程执行的操作,它实际上可能有助于调度程序避免死锁。
关于您在问题中给出的示例:在 CUDA 中,明确禁止将 __syncthreads()
放在 SIMD 发散代码中。编译器不会捕捉到它,但手册说“未定义的行为”。实际上,在 Fermi 之前的设备上,所有 __syncthreads()
都被视为相同的障碍。有了这个假设,您的代码实际上会终止而不会出现错误。不应该不依赖这种行为。
关于synchronization - CUDA/OpenCL 中的现实死锁示例,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/6426793/