什麼樣的變量消耗CUDA中的寄存器？

__global__ void add(int *c, const int* a, const int* b) 
{ 
    int x = blockIdx.x; 
    int y = blockIdx.y; 
    int offset = x + y * gridDim.x; 
    c[offset] = a[offset] + b[offset]; 
}

在上述例子中，我想x，y，offset被保存在寄存器中而什麼樣的變量消耗CUDA中的寄存器？

NVCC -Xptxas -v給出4 registers, 24+16 bytes smem
分析器顯示4個寄存器

和頭部PTX文件：

.reg .u16 %rh<4>; 
.reg .u32 %r<9>;  
.reg .u64 %rd<10>; 
.loc 15 21 0 

$LDWbegin__Z3addPiPKiS1_: 
.loc 15 26 0

任何人都可以澄清寄存器的使用情況如何？在Fermi中，每個線程的最大寄存器數爲63。在我的程序中，我想測試內核消耗太多寄存器的情況（因此變量可能必須自動存儲在本地內存中，從而導致性能下降）。然後在這一點上，我可以將一個內核分成兩個，這樣每個線程都有足夠的寄存器。假設SM資源足夠用於併發內核。

我不知道我是否正確。

來源

2012-07-14 user1525320

是你的問題「爲什麼這個代碼使用4個寄存器而不是3個？」如果是這樣，答案是這樣的：爲了添加'a [offset]'和'b [offset]'，必須獲取這兩個值。它必須存儲它在第一次獲取到的第一個地方，而它正在獲取另一個地方。所以還需要一個寄存器。 – 2012-07-14 12:20:43

謝謝您的回答，那麼我們可以說中間變量將被保存在寄存器中嗎？ – user1525320 2012-07-14 12:37:46

必要時，是的。在需要的時候並不總是很容易辨別，甚至可能因硬件目標而異。 – 2012-07-14 12:38:15

PTX中的寄存器分配與內核的最終寄存器消耗完全無關。 PTX僅是最終機器代碼的中間代表，並使用static single assignment form，這意味着PTX中的每個寄存器僅使用一次。一塊帶有數百個寄存器的PTX可以編譯成只有少數寄存器的內核。

寄存器分配由ptxas完成，作爲完全獨立的編譯過程（驅動程序靜態地或即時地執行），並且可以對輸入PTX執行大量代碼重新排序和優化以提高吞吐量並保存寄存器，這意味着PTX中的原始C或寄存器中的變量與組裝好的內核的最終寄存器計數之間幾乎沒有關係。

nvcc確實提供了一些方法來影響彙編程序的寄存器分配行爲。您有__launch_bounds__向編譯器提供了啓發式提示，這會影響寄存器分配，編譯器/彙編程序採用-maxrregcount參數（可能會導致寄存器溢出到本地內存，從而降低性能）。 volatile關鍵字用於對舊版本的基於nvopen64的編譯器產生影響，並可能影響本地內存溢出行爲。但是，您不能任意控制或引導原始C代碼或PTX彙編語言代碼中的寄存器分配。

來源

2012-07-14 13:15:23 talonmies

非常感謝，talonmies。所以我想我們的內核中沒有任何關於寄存器使用控制的事情？編譯器總是做很多事情。 – user1525320 2012-07-14 13:38:41

你有'__launch_bounds__'向編譯器提供啓發式提示，它可以影響寄存器分配，編譯器/彙編器則採用'-maxrregcount'參數。 'volatile'關鍵字用於與舊版本的nvopen64編譯器有所不同，並可能影響本地內存溢出行爲。但是你不能任意控制或引導原始C代碼中的寄存器分配。 – talonmies 2012-07-14 14:23:34

這有助於很多！再次感謝男士。 – user1525320 2012-07-14 14:48:29

什麼樣的變量消耗CUDA中的寄存器？

回答

相關問題