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()以前は1つのスレッドによってのみ到達は、いくつかの機能の中に隠さもちろんあったので、それを見つけるのは大変でした。私のセットアップ（sm50、gtx 980）では、このカーネルは実行され（宣言されたデッドロックはありません...）、アサーションは最初のワープの外側にあるすべてのスレッドで失敗します。

---

TL; DR

は、ブロック内のすべてのスレッドで呼び出されていない__syncthreads()を検出するための任意の標準的な方法はありますか？たぶん私は行方不明のいくつかのデバッガ設定？

アトミックとグローバルメモリを使用して状況を検出できる独自の（非常に遅い）checked__syncthreads()を作成できましたが、標準的な解決方法があります。

Answer 1

元のコードには、スレッド化されたデータ競合状態があります。
スレッド0は "p = true"に進んで実行することができますが、そのあと異なるスレッドがまだ進行していない可能性があり、結果を上書きするp = false行に戻ります。

この具体例

最も簡単な修正は単にPへの唯一のスレッド0のライトを持っている、

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