問題タブ [dynamic-parallelism]

dynamic parallelismこの新技術の説明から、ケプラーに関する情報はほとんどありませんが、同じワープでのスレッド制御フローの発散の問題が解決されたことを意味しますか？

デバイスrecursionコードからカーネルをランチすることができますが、異なるスレッドの制御パスを同時に実行できるということですか？

動的並列処理を使用して、CUDA のアルゴリズムを改善しようとしています。私の元の CUDA ソリューションでは、すべてのスレッドが各ブロックに共通の数値を計算します。私がやりたいことは、スレッドが共通の値を 1 回だけ計算する粗い (または低解像度) カーネルを最初に起動することです (すべてのスレッドが 1 つのブロックを表す場合のように)。次に、各スレッドは 1 ブロック (16x16 スレッド) の小さなグリッドを作成し、共通の値を渡す子カーネルを起動します。理論的には、多くの冗長な操作を節約できるため、高速になるはずです。しかし、実際には、ソリューションの動作が非常に遅く、その理由はわかりません。

これは非常に単純化されたコードであり、単なるアイデアです。

child_kernels の量は多く、スレッドごとに 1 つあり、約 400x400 スレッドが必要です。私が理解していることから、GPU はこれらすべてのカーネルを並行して処理する必要がありますよね?

または、子カーネルはどういうわけか順番に処理されますか?

私の結果は、パフォーマンスが元のソリューションよりも 10 倍以上遅いことを示しています。

頻繁な補間を必要とするいくつかのタスクを実行する、既に並列化された CUDA カーネルがあります。

だからカーネルがあります

この補間デバイス関数を 1 回以上呼び出します。

内挿アルゴリズムは、3 次元にわたって連続して WENO 内挿を行います。これは非常に並列化可能なタスクであり、早急に並列化したいと考えています!

構文complexStuff()を使用してホスト コードからカーネルを呼び出すことで、カーネルを簡単に並列化できることは明らかです。がすでに並列化<<<...>>>されていることも重要です。complexStuff()

しかし、何かを並列化する方法/CUDAデバイス関数内から新しいスレッドを作成する方法は明確ではありません...これは可能ですか? 誰か知っていますか？

CUDA プログラミング ガイドのセクション C.4.3.1.2 の下。「ネスティングと同期の深さ」には、次のように記載されています。

「親カーネルが cudaDeviceSynchronize() を呼び出さない場合に、親の状態用にスペースを予約する必要がないことをシステムが検出した場合、最適化が許可されます。この場合、明示的な親/子同期が発生しないため、プログラムは保守的な最大値よりもはるかに少なくなります。そのようなプログラムは、バッキング ストアの過剰な割り当てを避けるために、より浅い最大同期深度を指定できます。"

これは、コンパイラが動的並列処理で末尾再帰をサポートしていることを意味しますか? たとえば、自分自身を再帰的に呼び出すカーネルがあるとします。

ホストで起動:

末尾再帰をコンパイラで最適化できる場合、「親/子の同期は発生しない」ため、最大再帰レベル 24 に制限されますか? 制限されていない場合、コンパイラはどのように最適化を有効にできますか?

ありがとう！