Поддерживает ли nvcc оптимизацию хвостового вызова в динамическом параллелизме?
В разделе C.4.3.1.2 Руководства по программированию CUDA. "Глубина вложения и синхронизации", упоминается:
"Оптимизация разрешается, когда система обнаруживает, что ей не нужно резервировать место для состояния родителя в тех случаях, когда родительское ядро никогда не вызывает 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, поскольку "синхронизация между родителем и потомком никогда не происходит"? Если оно не ограничено, как компилятор может включить оптимизацию?
Спасибо!