CUDA/OpenCLの現実的なデッドロックの例

Question

私が書いているチュートリアルでは、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]; 

、それは1つがためのクリティカルセクションを実装しようとすると、初心者のCUDAプログラマでキャッチする、実際には簡単ですシンプルなデッドロックが適切CUDA CやOpenCLのC.

Answer 1

でもありませんすべてのスレッドによって最終的に実行される単一のスレッドです。それはこのように、多かれ少なかれ行く：

__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命令は、1つのスレッドがクリティカルセクションを終了すると、すべての他の人が1を取得しながら、それがセマフォを設定し、exacltyつのスレッドが前のために0に割り当てられますことを保証します他のスレッドがクリティカルセクションに入る可能性があるように0に戻します。

問題は、1つのスレッドがprev = 0になる間に、同じSIMDユニットに属する31のスレッドが値1を取得することです。if文では、CUDAスケジューラはその単一スレッドを保留（マスク）他の31スレッドに作業を継続させます。通常の状況では良い戦略ですが、この特定のケースでは決して実行されない1つのクリティカルセクションスレッドと31スレッドが無限を待っています。デッドロック。

また、whileループ外の制御フローを導くbreakが存在することにも注意してください。ブレーク命令を除外し、すべてのスレッドによって実行されるはずのifブロックの後にさらに操作がある場合、実際にスケジューラがデッドロックを回避するのに役立ちます。

質問に与えられた例について：CUDAでは、__syncthreads()をSIMD分岐コードに入れることは明示的に禁止されています。コンパイラはそれをキャッチしませんが、マニュアルは "未定義の動作"について言います。実際には、フェルミ・プリ・デバイス上では、すべて__syncthreads()が同じ障壁とみなされます。その前提で、あなたのコードは実際にエラーなく終了します。しかし、ではなく、はこの動作に依存しています。