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部分。