CUDAの__syncthreads（）と再帰

Question

私はこの__syncthreadsを行う方法

__device__ void foo(int k) { 
    if (some_condition) { 
    for (int i=0;i<8;i++) { 
     foo(i+k); // foo might take longer with some inputs 
     __syncthreads(); 
    } 
    } 
} 

のような再帰に__syncthreads（）を使用したい

（）今適用されますか？私はそれがブロック内でのみ適用されることを知っています。私が理解している限り、これは再帰の深さとは関係なくすべてのローカルスレッドに適用されます。しかし、もしこれを__syncthreads（）が特定の再帰深度に確実にしたいのであれば？それも可能ですか？私は再帰の深さを確認することができますが、どちらもうまくいかないと私は信じています。

代替手段がありますか？

私はCUDAデバイス> = 2.0

int __syncthreads_count(int predicate); 
int __syncthreads_and(int predicate); 
int __syncthreads_or(int predicate); 

のための3つのsyncthread拡張があることを見てきました。しかし、私は、彼らは、原子カウンターのように見えるので、彼らは役立つとは思いません。

Answer 1

ご存知のように、__syncthreads()は、ブロック内のすべてのスレッドがバリアに到達する場合にのみ安全です。つまり、条件内から__syncthreads()を呼び出す場合、条件はブロック内のすべてのスレッドで同じに評価されなければなりません。

再帰内で__syncthreads()の場合、ブロック内のすべてのスレッドが同じ深さまで再帰を実行する必要があります。そうでない場合、すべてのスレッドが同じバリアに到達するわけではありません。

Answer 2

代替手段はありますか？

はい、）（あなたが__syncthreadsについて言ったオフコース

Answer 3

あなたの機能ロジックを表現するために再帰パラダイムを使用していない事実であるブロック内のローカルのスレッドがゆえ、あなたが制御できないため、それだけで動作します他のブロックで何が起こっているか。削減のための最善の方法は、最初にブロック全体の大きさに等しい配列全体の縮小を行います。次に、配列をホストにコピーしないで、前の呼び出しのブロック数と同様の1ブロックとスレッドを持ち、後でサイズ1の配列をDeviceからHostにコピーする別の縮小を呼び出しません。最初の削減が生成されない限り、2回のコールcozの間でcudaThreadSynchronize（）を必ず使用してください。これは2段階の削減ですが、それは私のために働く。

乾杯！ saif