cuda SM寄存器限制

Question

我知道在一个SM上运行的块的数量受块号，线程，共享内存和寄存器的限制。是否有避免拥有太多寄存器的策略？我的意思是我只是不想太多，最终限制了我在一个SM上运行的块的数量。

Answer 1

寄存器数量的一个主要驱动因素是你在内核中声明的本地数据量。但是，PTX汇编器在重新使用寄存器方面可以做得很好，因此从PTX代码中计算出使用的数量并不总是容易 - 您需要运行ptxas以获得真正的答案。

Answer 2

用nvcc -Xptxas -v编译将打印出所提到的诊断信息Edric。此外，您可以使用__launch_bounds__限定符强制编译器保存寄存器。例如

__global__ void 
__launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor) 
MyKernel(...) 
{ 
    ... 
} 

保证大小maxThreadsPerBlock的至少minBlocksPerMultiprocessor块将适合在单个SM。有关__launch_bounds__的完整说明，请参阅CUDA Programming Guide的B.16部分。