如何在CUDA上实现深度递归

Question

在CUDA(数千个级别)上实现深度递归的最有效方法是什么？如果递归用于遍历树状数据结构,在哪里查找代码示例？

我刚刚使用Cuda Dynamic Parallelism在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 - 将评估每个节点是否需要由单独的线程处理,或者可能是整个warp或块甚至整个网格都可以分配给它.根据这一点,stack应该位于本地,共享或全局内存空间中,并且其成员函数应该在相应的执行单元(线程/块/网格)上统一运行.

请注意,如果您想stack在本地内存中使用每个线程,您将使用大量内存(10000个线程x 1000最大深度递归),并且您可能会遇到很多线程分歧,从而降低性能.

另一方面---每块stack需要更少的内存但__syncthreads()需要.

如果每个节点有足够的并行工作,我强烈建议对节点进行每个warp或每个块的处理.

最后,如果你在共享内存已经堆栈,但你发现你需要每经工作,你可以考虑使用的原子操作,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));
  } 
}

假设每个节点都有一个指向其父节点的指针(或者你可以在StackEntrystruct中引入这样的指针),你可以在树中传递参数.

但请注意,这会引入堆栈中条目之间的依赖关系.只要只有一个执行单元(线程/块/网格)从堆栈中推送/弹出,这就没问题了.但是,如果许多执行程序共享一个堆栈,使用前面讨论的工作窃取算法,它可能会破坏依赖关系.必须另外考虑以防止这种情况.

您可能希望重新组织确切StackEntry存储的内容以及何时将元素推入堆栈.以上方法不是唯一的方法!