问在CUDA中限制寄存器的使用：__launch_bounds__与maxrregcount

EN

这个问题的序言是引用CUDA C Programming Guide的话，

内核使用的寄存器越少，在多处理器上驻留的线程和线程块就越多，这可以提高性能。

现在，__launch_bounds__和maxregcount通过两种不同的机制限制寄存器的使用。

__launch_bounds__

nvcc通过平衡内核启动设置的性能和通用性来决定__global__函数要使用的寄存器数量。不同的说法是，选择使用寄存器的数量“保证”每个块的不同线程数和每个多处理器的块数。但是，如果在编译时可以得到每个块的最大线程数以及(可能)每个多处理器的最小块数的大致概念，那么可以使用这些信息来优化内核。换句话说

#define MAX_THREADS_PER_BLOCK 256
#define MIN_BLOCKS_PER_MP     2

__global__ void
__launch_bounds__(MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_MP)
fooKernel(int *inArr, int *outArr)
{
    // ... Computation of kernel
}

通知编译器可能的启动配置，以便nvcc能够以“最优”的方式为这种启动配置选择寄存器数量。

MAX_THREADS_PER_BLOCK参数是强制性的，而MIN_BLOCKS_PER_MP参数是可选的。还请注意，如果内核启动时每个块的线程数大于MAX_THREADS_PER_BLOCK，则内核启动将失败。

Programming Guide中对限制机制的描述如下：

如果指定了启动边界，编译器首先从它们派生内核应该使用的寄存器数量的上限L，以确保maxThreadsPerBlock线程的minBlocksPerMultiprocessor块(如果没有指定minBlocksPerMultiprocessor )可以驻留在多处理器上。然后，编译器以下列方式优化寄存器的使用：

• 如果初始寄存器使用率高于L，则编译器将其进一步减少，直至其小于或等于L，通常以牺牲更多的本地内存使用量和/或更多的指令数量为代价；

因此，__launch_bounds__可能导致注册溢出。

maxrregcount

maxrregcount是一个编译器标志，它通过强制编译器重新安排寄存器的使用，将雇用寄存器的数量严格限制在用户设置的数字上，与__launch_bounds__不一致。当编译器不能保持在规定的限制之下时，它只会将其泄漏到本地内存，这实际上是DRAM。即使这个局部变量存储在全局DRAM内存变量中，也可以缓存在L1，L2中。