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 节< /a> 的完整解释 __launch_bounds__。

Answer 2

寄存器数量的主要驱动因素之一是您在内核中声明的本地数据量。然而，PTX 汇编器在重用寄存器方面可以做得相当好，因此从 PTX 代码中计算出将使用多少寄存器并不总是那么容易 - 您需要运行 ptxas 来获取真正的答案。