CUDA/OpenCL中的逼真死锁示例

对于我正在编写的教程，我正在寻找一个“现实”且简单的由SIMT/SIMD无知造成的死锁例子。CUDA/OpenCL中的逼真死锁示例

我想出了这个片段，这似乎是一个很好的例子。

任何输入，将不胜感激。

… 
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.

来源

2011-06-21 Framester

作为一个“现实”的例子，这似乎太复杂了。我只会在条件中使用'get_local_id（0）>常量'，并用注释替换“业务代码”（赋值）'/ *做一些事情* /'和'/ *做另一件事* /'。尽管如此，我认为StackOverflow并不是讨论的最佳地点，它是一个提问和回答的地方。 –

一个简单的僵局，实际上是容易被新手程序员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分配到上一个，而所有其他人得到1.当一个线程完成它的关键部分，它集信号回到0，以便其他线程有机会进入临界区。

问题是，当1个线程获得prev = 0时，属于相同SIMD单元的31个线程获得值1.在if语句中，CUDA调度程序将该单个线程置于保持状态（将其屏蔽）并让其他31线程继续他们的工作。在正常情况下，这是一个很好的策略，但在这种特殊情况下，您最终会得到1个永不执行的关键节线程和31个等待无穷大的线程。僵局。

另请注意，break的存在导致控制流程在while循环之外。如果你忽略了中断指令，并且在应该由所有线程执行的if块之后有更多的操作，它实际上可以帮助调度器避免死锁。

关于在问题中给出的示例：在CUDA中明确禁止将__syncthreads()放在SIMD分叉代码中。编译器不会捕获它，但手册中提到了“未定义的行为”。实际上，在费米器件前，所有的__syncthreads()都被视为同样的障碍。有了这个假设，你的代码实际上会终止而没有错误。一个应该不是虽然依赖于此行为。

来源

2011-06-21 19:22:33 CygnusX1

CUDA/OpenCL中的逼真死锁示例

回答

相关问题