一个CUDA程序如果使用的寄存器数量过多,会导致在SM上同时驻留的线程和block数量减少,继而导致程序性能不足。
__launch_bounds__
和maxrregcount
都可以用来限制cuda程序的寄存器数量,但是两者是不同的机制。
__launch_bounds__
__global__ void
__launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor)
MyKernel(...)
{
...
}
- maxThreadsPerBlock 指定了每个block最多可以包含多少个线程。它对应了PTX中的
.maxntid
指令。 - minBlocksPerMultiprocessor 是可选的,它指定了每个SM上最少需要驻留多少block,它对应了PTX的
.minnctapersm
指令。
指定这两个参数以后,compiler会推出每个线程的寄存器使用上限值L。
compiler将通过下面的方法来优化寄存器使用:
- 如果compiler第一次得出的寄存器使用量高于L,那么它将通过spill寄存器内容到local memory和增加指令数量(通过更多的寄存器内容交换)来减少寄存器使用量,直到小于等于L。
- 如果第一次得出的寄存器量低于L,
-
- 如果指定了maxThreadsPerBlock但是没有指定minBlocksPerMultiprocessor,compiler则会尽量减少寄存器以便满足这个max。
- 如果两个都指定了,则compiler会尽量增加寄存器数量(直到L)来减少指令数量隐藏指令延迟。
maxrregcount
maxrregcount是一个编译参数,简单暴力的将所有的kernel函数寄存器(除了使用launch_bound限制的kernel函数以外)使用的最大量限制为对应的数值。
它只关心寄存器的使用数量,我感觉它的优化策略,跟launch_bound第一条差不多。
- 如果原来default编译出来寄存器数量超过了maxrregcount,就尽量减少寄存器数量。
- 但是如果原来的寄存器数量很少,而maxrregcount限制值较大,寄存器数量可能减少、不变或者增加
Reference
Limiting register usage in CUDA: launch_bounds vs maxrregcount
CUDA Programming Guide
Comments