cuda SM注册限制

Question

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

Answer 1

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

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

保证至少minBlocksPerMultiprocessor块大小maxThreadsPerBlock适合单个SM.有关完整说明,请参见" CUDA编程指南"的 B.16节__launch_bounds__.