在CUDA中，如何檢測__syncthreads（）未被塊中的所有線程調用？

Question

我剛剛遇到了一個奇怪的問題，很難在CUDA中產生原來涉及未定義行爲的問題。我想讓線程0在所有線程應該使用的共享內存中設置一些值。

__shared__ bool p; 
p = false; 
if (threadIdx.x == 0) p = true; 
__syncthreads(); 
assert(p); 

現在assert(p);隨機看似失敗，因爲我周圍鏟的代碼，並評論它出去找問題。

我曾使用過這種結構中有效以下未定義行爲方面：

#include <assert.h> 

__global__ void test() { 
    if (threadIdx.x == 0) __syncthreads(); // call __syncthreads in thread 0 only: this is a very bad idea 
    // everything below may exhibit undefined behaviour 


    // If the above __syncthreads runs only in thread 0, this will fail for all threads not in the first warp 
    __shared__ bool p; 
    p = false; 
    if (threadIdx.x == 0) p = true; 
    __syncthreads(); 
    assert(p); 
} 

int main() { 
    test << <1, 32 + 1 >> >(); // nothing happens if you have only one warp, so we use one more thread 
    cudaDeviceSynchronize(); 
    return 0; 
} 

早期__synchthreads()只達到由一個線程是進程隱藏在一些功能上的，所以這是很難找到。在我的設置（sm50，gtx 980）上，這個內核運行完畢（沒有像公佈的那樣死鎖......），並且第一個warp之外的所有線程斷言失敗。

---

TL; DR

是否有檢測__syncthreads()不被塊中的所有線程調用任何標準的方式？也許我錯過了一些調試器設置？

我可以構建自己的（非常慢）checked__syncthreads()，可以使用可能原子和全局內存檢測情況，但我寧願有一個標準的解決方案。

Answer 1

您的原始代碼中存在線程數據競爭條件。
線程0可能前進到並執行「p = true」，但在此之後，不同的線程可能根本沒有進展，並且仍然會回到p = false行，覆蓋結果。

對於這個具體的例子最簡單的解決將簡單地僅具有線程0寫入至p，像

__shared__ bool p; 
if (threadIdx.x == 0) p = true; 
__syncthreads(); 
assert(p);