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()必要になります。

ノードごとに十分な並列作業が行われる場合は、ノードのワープごとまたはブロックごとの処理を強くお勧めします。

最後に、共有メモリにスタックがあり、ワープごとに作業が必要であることがわかった場合は、アトミック操作を使用することを検討し、ワープ間で作業のバランスを取るためにワークスティーリング手法を導入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構造体にそのようなポインターを導入できる) と仮定すると、パラメーターをツリーに渡すことができます。

ただし、これによりスタック内のエントリ間に依存関係が生じることに注意してください。これは、1 つの実行ユニット (スレッド/ブロック/グリッド) だけがスタックからプッシュ/ポップしている限り問題ありません。ただし、1 つのスタックが多くのエグゼキューターによって共有されている場合、前述のワーク スティーリング アルゴリズムを使用すると、依存関係が壊れる可能性があります。それを防ぐには、追加の検討が必要です。

正確に何StackEntryが格納されているか、要素がいつスタックにプッシュされるかを再編成したい場合があります。上記のアプローチだけではありません。

于 2013-01-14T01:11:15.803 に答える