1

CUDA 5 プログラミング ガイドでは、次のように述べられています。

起動は 24 世代の深さまで継続する可能性がありますが、この深さは通常、GPU で使用可能なリソースによって制限されます

私の質問は次のとおりです。

  • GPU 上の CUDA ランタイムは、常に 24 の深さを達成できることを保証し、場合によっては 24 を超えることさえありますか (ケース A)? それとも、24 が絶対的な最大制限であり、実行時にこの数に実際に到達しない可能性があることを意味しますか (ケース B)?

  • ケース B の場合、カーネルが GPU で起動され、十分なリソースがない場合はどうなりますか? 打ち上げ失敗？（こんなことになったらおかしい！）

CUDA プログラムを作成する予定で、Kepler アーキテクチャを活用したいと考えています。私のアルゴリズムでは、通常 15 ～ 19 のレベルで関数の再帰が絶対に必要です (再帰レベルはデータ構造にバインドされています)。

参照: TechBrief_Dynamic_Parallelism_in_CUDA_v2.pdf

4

1 に答える 1

1

CUDA は、再帰の深さ 1 が達成されることを保証しません。同様に、従来の OS は、新しいプロセス/スレッドの起動が成功することを保証しません。たとえば、次のプログラムがあるとします。

int main() {
    pid_t pid;
    while (pid = fork ());
    while (true) {
        dummy<<<1, 1024>>> ();
    }
}

__global void dummy() {}

ある時点で何かが失敗します - CPU または GPU メモリが不足します。同様に、GPU では失敗する可能性があります (エラーを返します - CUDA または fork のいずれかが -1 を返します)。

別の見方をすると、最悪の場合、各起動には(2^31-1)^2*(2^10-1) ≃ 2^72それぞれスレッドを持つブロックが存在する可能性があります。2^10つまり、単一の起動では、2^82スレッドを持つことができます。現在、各再帰は指数関数的であるため、最悪の場合、起動後にスレッドを終了しても、2^1968スレッドのスケジューリングを保証する必要があります。各スレッドの状態が 1/32 ビットの場合、ワープが終了したかどうかにかかわらず、2^1945GiB のメモリが必要になります (これは、観測された宇宙の情報容量よりも「わずかに」多く、つまり2^15952 倍以上です)。

したがって、それは間違いなくケース B であり、ケース A である可能性はありません (ワープの状態には、少なくとも命令ポインターが含まれている必要があります)。分岐要因に応じて、同期する場合、アルゴリズムの 15-19 の再帰の深さが達成可能になる場合があります。

編集:再帰的な起動の代わりに単純な再帰を意味する場合、実際にはスタックによって制限されます。正確なコードによっては、Fermi+ では事実上無限になる場合があります (Tesla 生成は再帰 IIRC をサポートしていません)。同様に、最小の深さが保証されているわけではありません - スタック/ローカルメモリに大きな配列を割り当てると、スペースが不足します (オプティマイザはそれをうまく取り除きます)。

于 2013-01-13T11:07:26.247 に答える