如何在CUDA上实现深度递归

Question

在CUDA (千层)上实现深度递归的最有效方法是什么?如果递归用于遍历类似树的数据结构，那么在哪里可以找到这方面的代码示例？

我刚刚使用Cuda动态并行实现了K20 GPU上的递归，但发现由于参数cudaLimitDevRuntimeSyncDepth，有24个级别的限制。

我想要达到最大。大数据的速度和缩放。

Answer 1

根据我的经验，在CUDA中管理递归的最可靠和最有效的方法是手动管理递归堆栈，并将功能“扁平化”。例如，如果您正在遍历一棵二叉树，它将如下所示：

while (!stack.isEmpty()) {
  Node n = stack.pop();
  ... //do stuff with n
  if (!n.isLeaf()) {
    stack.push(n.left());
    stack.push(n.right());
  } 
}

上述技术可以帮助任何代码(CUDA或单线程CPU).堆栈功能必须由您实现，因为您不想使用STL。

下一步--更具体到CUDA --将评估每个节点是否需要被单独的线程处理，或者整个扭曲或块，甚至整个网格都可以分配给它。根据这一点，stack应该位于本地、共享或全局内存空间中，其成员函数应该在相应的执行单元(线程/块/网格)中一致运行。

请注意，如果您想在本地内存中使用每线程stack，您将使用大量内存(10000线程x1000max深度递归)，并且可能会遇到大量线程差异，降低性能。

另一方面，每块stack将需要较少的内存，但需要__syncthreads()。

如果有足够的并行工作做每个节点，我强烈建议每翘曲或每块处理节点。

最后，如果您在共享内存中有堆栈，但发现您需要每次工作，您可以考虑为push和pop使用原子操作，并引入窃取工作的技术，以便在翘曲之间更好地平衡您的工作。如果您需要通过全局内存中的单个堆栈来处理每个节点的块，则还可以执行偷取工作。

编辑：--如果您需要在树上行走，在处理完树之后，您可以将向上的方向推到树中。

struct StackEntry {
    Node* node;
    bool goingUp;
};

while (!stack.isEmpty()) {
  StackEntry entry = stack.pop();
  ... //do stuff with entry.node
  if (!entry.goingUp && !entry.node->isLeaf()) {
    stack.push(StackEntry(entry.node->left(),false));
    stack.push(StackEntry(entry.node->right(),false));
    stack.push(StackEntry(entry.node,true));
  } 
}

假设每个节点都有一个指向其父节点的指针(或者您可以在StackEntry结构中引入这样的指针)，则可以将参数传递到树中。

但是，请注意，这会在堆栈中的条目之间引入依赖关系。只要只有一个执行单元(线程/块/网格)从堆栈中推送/弹出就可以了。但是，如果一个堆栈被多个执行者共享，并且使用前面讨论过的工作窃取算法，它可能会打破依赖关系。必须进一步考虑以防止这种情况发生。

您可能想要重新组织确切的StackEntry存储的内容以及元素何时被推到堆栈上。以上方法并不是唯一的方法！