Информации по dynamic parallelism Кеплера очень мало, из описания этой новой технологии, значит решается вопрос расхождения потока управления потоками в одном и том же варпе?
Это позволяет recursion и запускать ядро из кода устройства, означает ли это, что путь управления в разных потоках может выполняться одновременно?
Взгляните на этот документ.
Динамический параллелизм, дивергенция потока и рекурсия — разные понятия. Динамический параллелизм — это возможность запуска потоков внутри потока. Это означает, например, что вы можете сделать это
__global__ void t_father(...) {
...
t_child<<< BLOCKS, THREADS>>>();
...
}
Я лично исследовал эту область, когда вы делаете что-то подобное, когда t_father запускает t_child, все ресурсы vga снова распределяются между ними, и t_father ждет, пока все t_child не закончатся, прежде чем он сможет продолжить (посмотрите также это статья Слайд 25)
Рекурсия доступна со времен Fermi и представляет собой возможность для потока вызывать себя без какой-либо другой реконфигурации потока/блока.
Что касается расхождения потока, я думаю, мы никогда не увидим поток внутри варпа, одновременно выполняющий разные коды.
Нет. Концепция деформации все еще существует. Все потоки в варпе являются SIMD (Single Instruction Multiple Data), что означает, что они одновременно выполняют одну инструкцию. Даже когда вы вызываете дочернее ядро, GPU назначает один или несколько деформаций вашему вызову. При использовании динамического параллелизма помните о трех вещах:
Максимально возможное значение — 24 (CC = 3,5).
Количество одновременно работающих динамических ядер ограничено (по умолчанию 4096), но может быть увеличено.
Держите родительское ядро занятым после вызова дочернего ядра, в противном случае есть большая вероятность, что вы потратите ресурсы впустую.
В этой презентации NVidia есть пример исходного кода cuda. на слайде 9.
__global__ void convolution(int x[])
{
for j = 1 to x[blockIdx]
kernel<<< ... >>>(blockIdx, j)
}
Далее показано, как часть управляющего кода CUDA перемещается в графический процессор, чтобы ядро могло порождать другие функции ядра в частичных доменах домена разного размера (слайд 14).
Глобальный вычислительный домен и его разбиение на разделы по-прежнему статичны, поэтому вы не можете пойти и изменить это вычисление ВО ВРЕМЯ GPU, например, на порождать больше запусков ядра, потому что вы еще не достигли конца своей функции оценки. Вместо этого вы предоставляете массив, содержащий количество потоков, которые вы хотите создать с определенным ядром.
