cuda SM寄存器限制
问题描述:
我知道在一个SM上运行的块的数量受块号,线程,共享内存和寄存器的限制。是否有避免拥有太多寄存器的策略?我的意思是我只是不想太多,最终限制了我在一个SM上运行的块的数量。cuda SM寄存器限制
答
寄存器数量的一个主要驱动因素是你在内核中声明的本地数据量。但是,PTX汇编器在重新使用寄存器方面可以做得很好,因此从PTX代码中计算出使用的数量并不总是容易 - 您需要运行ptxas
以获得真正的答案。
答
用nvcc -Xptxas -v
编译将打印出所提到的诊断信息Edric。此外,您可以使用__launch_bounds__
限定符强制编译器保存寄存器。例如
__global__ void
__launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor)
MyKernel(...)
{
...
}
保证大小maxThreadsPerBlock
的至少minBlocksPerMultiprocessor
块将适合在单个SM。有关__launch_bounds__
的完整说明,请参阅CUDA Programming Guide的B.16部分。