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中的逼真死鎖示例

回答

相關問題