調用__device__函數是否會影響CUDA中使用的寄存器數量？

Question

我已經閱讀過各種地方，__device__函數幾乎總是由CUDA編譯器內聯。那麼說，當我將代碼從內核移動到由內核調用的__device__函數中時，所使用的寄存器數量（通常）不會增加？

作爲一個例子，下面的代碼段使用相同數量的寄存器嗎？他們有同樣的效率嗎？

SNIPPET 1

__global__ void manuallyInlined(float *A,float *B,float *C,float *D,float *E) { 
    // code that manipulates A,B,C,D and E 
} 

SNIPPET 2

__device__ void fn(float *A,float *B,float *C,float *D,float *E) { 
    // code that manipulates A,B,C,D and E 
} 


__global__ void manuallyInlined(float *A,float *B,float *C,float *D,float *E) { 
    fn(A,B,C,D,E); 
} 

Answer 1

最終答案只能通過使用工具來確定（編譯-Xptxas -v，或者使用分析器中的一個），但一般的答案是調用__device__函數可以影響所用寄存器的數量（如w性能和效率）。

根據您的文件組織，以及如何編譯代碼，__device__功能可能是inlined。如果內聯，通常會給優化編譯器（ptxas，主要）提供適應註冊使用情況的最佳機會。 （請注意，至少在理論上，這種「適應」可能導致或多或少的寄存器被使用，然而，內聯的情況通常導致編譯器同時使用更少的寄存器和更高的性能，但編譯器主要是優化了更高的性能，而不是更少的寄存器使用。）

另一方面，如果它沒有內聯，那麼它必須作爲一個普通的函數調用來處理。像許多其他計算機體系結構一樣，函數調用包括設置堆棧幀以傳遞變量，然後將控制權轉移給函數。在這種情況下，編譯器是更受限制的，因爲：

1. 它必須從堆棧框架移動由所述函數中使用的變量來/
2. 基於「周圍」的代碼，不能進行其他優化，因爲它不知道周圍的代碼是什麼。編譯器必須以獨立的方式處理__device__函數。

所以，如果函數可以內聯，你的兩種方法應該沒有太大的區別。如果函數不能內聯，那麼在上述兩種方法中，寄存器使用通常會有明顯的差異。

可能影響編譯器是否將嘗試內聯__device__功能的一些明顯的因素是：

1. 如果__device__功能是在一個單獨的編譯單元從調用它的__global__或其他__device__功能。在這種情況下，唯一可行的方法是通過CUDA separate compilation and linking，也稱爲設備鏈接。在這種情況下，編譯器不會（不能）內聯該函數。

2. 如果指定__noinline__compiler directive。請注意，這只是編譯器的一個提示;它可能被忽略。