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 Device> = 2.0有3个syncthread扩展

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

但我不认为他们会帮助他们,因为他们看起来像一个原子计数器.

Answer 1

如您所知,__syncthreads()只有在块内的所有线程到达屏障时才是安全的.这意味着如果您__syncthreads()在条件内调用,则条件必须在块中的所有线程上求值相同.

对于__syncthreads()递归内,这意味着块中的所有线程必须执行到相同深度的递归,否则并非所有线程都将到达相同的屏障.