cuda SM предел регистрации

Question

Я знаю, что количество блоков, работающих на одном SM, ограничено номером блока, потоками, общей памятью и регистром.Есть ли стратегия, позволяющая избежать слишком большого количества регистров?Я имею в виду, что я просто не хочу иметь их слишком много, в конечном итоге это ограничивает количество блоков, которые я запускаю на одном SM.

Answer 1

Компиляция с nvcc -Xptxas -v распечатает диагностическую информацию, упомянутую Эдриком.Кроме того, вы можете заставить компилятор сохранять регистры, используя квалификатор __launch_bounds__.Например,

__global__ void
__launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor)
MyKernel(...)
{ 
   ...
}

гарантирует, что по крайней мере minBlocksPerMultiprocessor блоков размером maxThreadsPerBlock поместятся на одном SM.См. Раздел B.16 Руководство по программированию CUDA для полного объяснения __launch_bounds__.

Answer 2

Одним из основных драйверов для количества регистров является количество локальных данных, которые вы объявляете в своем ядре.Тем не менее, ассемблер PTX может неплохо выполнить повторное использование регистров, поэтому не всегда легко определить, сколько из них будет использовано из кода PTX - для получения реального ответа необходимо запустить ptxas.*