CUDA プログラミング ガイドのセクション C.4.3.1.2 の下。「ネスティングと同期の深さ」には、次のように記載されています。
「親カーネルが cudaDeviceSynchronize() を呼び出さない場合に、親の状態用にスペースを予約する必要がないことをシステムが検出した場合、最適化が許可されます。この場合、明示的な親/子同期が発生しないため、プログラムは保守的な最大値よりもはるかに少なくなります。そのようなプログラムは、バッキング ストアの過剰な割り当てを避けるために、より浅い最大同期深度を指定できます。"
これは、コンパイラが動的並列処理で末尾再帰をサポートしていることを意味しますか? たとえば、自分自身を再帰的に呼び出すカーネルがあるとします。
__global__ void kernel(int layer){
if(layer>65535){
return;
}
printf("layer=%d\n",layer);
kernel<<<1,1>>>(layer+1);
}
ホストで起動:
kernel<<<1,1>>>(0);
末尾再帰をコンパイラで最適化できる場合、「親/子の同期は発生しない」ため、最大再帰レベル 24 に制限されますか? 制限されていない場合、コンパイラはどのように最適化を有効にできますか?
ありがとう!