Я знаю, что число блоков, работающих на одном SM, ограничено номером блока, потоками, общей памятью и регистром. Есть ли стратегия избежать слишком большого количества регистров? Я имею в виду, что я просто не хочу иметь слишком много из них, в конечном итоге он ограничивает количество блоков, которые я запускаю на одном SM.Ограничение регистра Cuda SM
Одним из основных драйверов для количества регистров является количество локальных данных, которые вы заявляете в своем ядре. Тем не менее, ассемблер PTX может отлично справиться с повторным использованием регистров, поэтому не всегда легко определить, сколько будет использовано из кода PTX - вам нужно запустить ptxas, чтобы получить реальный ответ.
Компиляция с nvcc -Xptxas -v распечатает диагностическую информацию, о которой упоминал Эрик. Кроме того, вы можете заставить компилятор сохранять регистры с помощью квалификатора __launch_bounds__. Например
__global__ void
__launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor)
MyKernel(...)
{
...
}
гарантирует, что по крайней мере minBlocksPerMultiprocessor блоки размером maxThreadsPerBlock умещаются на одном СМ. См. Раздел B.16 of CUDA Programming Guide для полного объяснения __launch_bounds__.