CUDA __syncthreads() 和递归答案

【问题标题】：CUDA __syncthreads() and recursionCUDA __syncthreads() 和递归
【发布时间】：2011-07-19 10:01:06
【问题描述】：

我想使用 __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 个同步线程扩展

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

但我认为它们不会有帮助，因为它们看起来像一个原子计数器。

【问题讨论】：

我没有确切的答案，因为我自己从未做过类似的事情，只是为了检查您是否知道，您在代码中输入的 some_condition 必须对相同的块，否则会死锁。
是的，这也是我害怕的。
请问您能澄清您的问题吗？我真的不明白你在这里问什么。

标签： recursion cuda

【解决方案1】：

如您所知，__syncthreads() 只有在块内的所有线程都到达屏障时才是安全的。这意味着，如果您在条件内调用 __syncthreads()，则该条件在块内的所有线程上的计算结果必须相同。

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

【讨论】：

您的推理是有道理的，但我可以想象，由于递归需要 fermi GPU，因此递归深度无关紧要，只需在代码中的位置即可。它可能可以查看堆栈深度，是的，但是为什么，这会引入大量潜在问题（如死锁）。我试图找到有关此的更多信息。这是在某处定义的吗？最简单且最可能的解决方案是：不要在递归中使用它
我会更强烈地说：除非必须，否则不要在 CUDA 中使用递归。每个线程都必须维护自己的堆栈，这会导致大量额外的片外内存访问，如果您可以用迭代替换递归，则不需要这些访问。如果不能，那么您也许可以在共享内存中维护一个更简单的堆栈。或者您可以在共享内存或寄存器中维护堆栈的顶部几层，从而减少片外访问的总数（通常用于 GPU 光线追踪）。至于syncthreads()，在任何非发散代码、递归或其他代码中使用都是安全的。
好的，感谢您对此的澄清。我认为 cuda 堆栈可能比自制的堆栈更有效。由于在此之前我没有太多接触过 cuda，这是表达它的最简单方法。无论如何，我目前正在重写所述代码（无递归和无堆栈），但我对递归中的 __syncthreads 行为非常感兴趣。因为我在启用 __syncthreads() 的情况下运行了这段代码，并且它可以正常工作并且令人惊讶地没有死锁。

【解决方案2】：

有没有可能的替代品？

是的，不要使用递归范式来表达你的函数逻辑

【讨论】：

【解决方案3】：

当然，您所说的 __syncthreads() 是正确的，它仅适用于块内的本地线程，因此您无法控制其他块中发生的事情。减少的最好方法是首先对整个数组进行减少，这将一般一个数组等于块的大小。然后不要将数组复制回主机，而是调用另一个缩减，这将有 1 个块和线程，类似于之前调用中的块数，然后将大小为 1 的数组从设备复制到主机。但请确保在两次调用之间使用 cudaThreadSynchronize() 除非生成第一个缩减，否则您可以进行缩减。这是减少了两步，但它对我有用。

干杯！！！赛夫

【讨论】：