限制CUDA中的寄存器使用:__ launch_bounds__ vs maxrregcount

Question

来自NVIDIA CUDA C编程指南:

可以使用maxrregcount编译器选项控制寄存器使用,也可以按启动边界中的描述启动边界.

从我的理解(并纠正我,如果我错了),虽然-maxrregcount限制整个.cu文件可能使用的寄存器数量,__launch_bounds__限定符定义maxThreadsPerBlock和minBlocksPerMultiprocessor每个__global__内核.这两个完成相同的任务,但有两种不同的方式.

我的用法要求我40每个线程都有寄存器以最大化性能.因此,我可以使用-maxrregcount 40.我也可以40通过使用强制寄存器,__launch_bounds__(256, 6)但这会导致加载和存储寄存器溢出.

导致这些寄存器泄漏的两者之间有什么区别？

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,则内核启动将失败.

限制机制Programming Guide如下所述:

如果指定了启动边界,则编译器首先从它们中获取L内核应该使用的寄存器数量的上限,以确保线程的minBlocksPerMultiprocessor块(或者如果minBlocksPerMultiprocessor未指定单个块) maxThreadsPerBlock可以驻留在多处理器上.然后编译器以下列方式优化寄存器使用:

• 如果初始寄存器使用率高于L,则编译器将其进一步减小,直到它变得小于或等于L,通常以更多本地存储器使用和/或更多指令数为代价;

因此,__launch_bounds__可导致注册溢出.

maxrregcount

maxrregcount是一个编译器标志__launch_bounds__,通过强制编译器重新安排其寄存器的使用,简单地将所使用的寄存器的数量限制为由用户设置的数字,与其不一致.当编译器不能保持低于强制限制时,它只会将其溢出到本地存储器,实际上DRAM.即使这个局部变量存储在全局DRAM内存变量中,也可以缓存在L1,L2中.