4

在 CUDA（数千级）上实现深度递归的最有效方法是什么，如果递归是为了遍历树状数据结构，在哪里可以找到代码示例？

我刚刚使用 Cuda Dynamic Parallelism 在 K20 GPU 上实现了递归，但发现由于参数 cudaLimitDevRuntimeSyncDepth 有 24 个级别的限制

我想达到最大值。大数据的速度和扩展。

4

1 回答 1

7

根据我的经验，在 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 个线程 x 1000 最大深度递归），并且您可能会遇到很多线程分歧，从而降低您的性能。

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

如果每个节点有足够的并行工作，我强烈建议对节点进行 per-warp 或 per-block 处理。

最后，如果您在共享内存中有堆栈，但您发现需要为每个 Warp 工作，您可以考虑使用原子操作 for pushandpop并引入工作窃取技术来更好地平衡您的工作在 Warp 之间。如果您需要通过在全局内存中使用单个堆栈来进行每个节点的块处理，也可以进行工作窃取。

编辑： 如果您需要向上走树，在将其向下处理后，您可以稍后将向上的方向推入树中。

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存储的内容以及将元素推入堆栈的时间。上述方法不是唯一的一种！

于 2013-01-14T01:11:15.803 回答