2

在 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 的限制，因为“父/子同步永远不会发生”？如果不限制，编译器怎么开启优化呢？

谢谢！

4

0 回答 0