如何在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線程x 1000最大深度遞歸），並且您可能會遇到太多線程分歧以降低性能。

另一方面---每塊stack將需要較少的內存，但__syncthreads()將需要。

如果每個節點有足夠的並行工作，我強烈建議節點的每個warp或per-block處理。

---

最後，如果你已經堆放在共享內存，但你發現你需要工作每經編，您可以考慮使用原子操作爲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正在存儲的內容以及何時將元素壓入堆棧。上述方法不是唯一的方法！