Поддерживает ли 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, поскольку "синхронизация между родителем и потомком никогда не происходит"? Если оно не ограничено, как компилятор может включить оптимизацию?

Спасибо!

0 ответов

Другие вопросы по тегам