限制寄存器使用在CUDA：__launch_bounds__ VS maxrregcount

Question

从NVIDIA CUDA C Programming Guide：

注册用法可以使用编译器maxrregcount 选项控制或如启动界描述启动边界。

从我的理解（和纠正我，如果我错了），而-maxrregcount限制寄存器整个.cu文件可以使用的数量，__launch_bounds__预选赛定义每个__global__内核maxThreadsPerBlock和minBlocksPerMultiprocessor。这两个人完成同样的任务，但有两种不同的方式。

我的用法要求我为每个线程配备40寄存器以最大限度地提高性能。因此，我可以使用-maxrregcount 40。我也可以使用__launch_bounds__(256, 6)强制40寄存器，但这会导致加载&存储寄存器溢出。

导致这些寄存器溢出的两者有什么区别？

Answer 1

这个问题的前言是，引述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，则内核启动将失败。

如果指定发射边界，编译器首先从它们 导出上限L上的内核应使用 确保寄存器的数目：

该限制机构设置在Programming Guide如下描述那的 maxThreadsPerBlock线程minBlocksPerMultiprocessor块（或单个块 如果未指定minBlocksPerMultiprocessor是）可以驻留在多处理器。所述 编译器将优化寄存器以下列方式使用：在

• 如果初始寄存器使用比L更高，编译器进一步降低它，直到它变得小于或等于L，通常更多的本地为代价内存使用量和/或更高数量的指令;

因此，__launch_bounds__可导致寄存器溢出。

maxrregcount

maxrregcount是简单地用__launch_bounds__ hardlimits使用的寄存器的数目由用户设定的数目，在方差，强制编译器将重新安排其使用的寄存器的编译器标志。当编译器不能保持低于规定的限制时，它只会将其泄漏到本地内存中，实际上它是DRAM。即使这个局部变量存储在全局变量DRAM中，内存变量也可以缓存在L1，L2中。