強制CUDA使用寄存器作爲變量

我在內核中有很多未使用的寄存器。我想告訴CUDA使用一些寄存器來保存一些數據，而不是每次需要時都讀取全局數據。（我不能夠使用共享MEM）強制CUDA使用寄存器作爲變量

__global__ void simple(float *gData) { 
float rData[1024]; 
for(int i=0; i<1024; i++) { 
    rData[i]=gData[i]; 
    } 
// work on the data here 
}

編譯瓦特/：NVCC -arch sm_20 --ptxas選項= -v simple.cu，我也得到
0字節堆棧幀，0字節溢出存儲，0字節溢出負載
使用2個寄存器，40個字節CMEM [0]

__global__ void simple(float *gData) { 
register float rData[1024]; 
for(int i=0; i<1024; i++) { 
    rData[i]=gData[i]; 
    } 
// work on the data here 
}

寄存器聲明什麼都不做。
0字節堆棧幀，0字節溢出存儲，0字節溢出負載
使用2個寄存器，40個字節CMEM [0]

__global__ void simple(float *gData) { 
volatile float rData[1024]; 
for(int i=0; i<1024; i++) { 
    rData[i]=gData[i]; 
    } 
// work on the data here 
}

揮發性聲明創建堆棧存儲：
4096字節堆棧幀， 0字節溢出店，二手0字節溢出負荷
21名抵抗者，40個字節CMEM [0]

1）是否有一個簡單的方法來告訴編譯器使用寄存器空間的變量？
2）'堆棧幀'在哪裏：寄存器，全局mem，本地mem，...？什麼是堆棧框架？（由於當沒有所述GPU具有堆疊的虛擬堆疊？）
3）simple.ptx文件基本上是空的：（NVCC -arch sm_20 -ptx simple.cu）

.loc 2 14 2 
ret;

任何想法，我可以找到真正的機器/編譯代碼？

來源

2012-08-28 Doug

編譯器優化了整個代碼，因爲它不修改任何非瞬態狀態。 – njuffa

每個線程要求1024個寄存器是一個非常高的順序。大多數內核每個線程需要數十個寄存器。如果你想確保編譯器可以使用一個寄存器作爲變量，它需要是一個標量（即不是你在'for'循環中索引的數組）。 –

在哪裏/什麼堆棧框架答案可以在這裏找到：http://stackoverflow.com/questions/7810740/where-does-cuda-allocate-the-stack-frame-for-kernels – Doug

動態索引數組無法存儲在寄存器中，因爲GPU寄存器文件不是動態可尋址的。
標量變量由編譯器自動存儲在寄存器中。
靜態索引（即，其中的索引可以在編譯時來確定）小陣列（比如，小於16個浮子）可以被存儲在由編譯器寄存器。

SM 2.0 GPU（費米）只支持每個線程多達63個寄存器。如果超過此值，寄存器值將從緩存層次結構支持的本地（片外）內存溢出/填充。 SM 3.5 GPU將其擴展到每個線程最多255個寄存器。一般來說，正如Jared所提到的，每個線程使用太多的寄存器是不期望的，因爲它減少了佔用，因此減少了內核中的延遲隱藏能力。 GPU在並行性方面蓬勃發展，並通過覆蓋來自其他線程的工作來延遲內存延遲。

因此，你可能不應該優化陣列到寄存器。相反，請確保跨線程訪問這些陣列的內存儘可能接近順序，以便最大程度地實現合併（即最大限度地減少內存事務）。

你給可以是用於共享存儲器如果的情況下的例子：

在塊許多線程使用相同的數據，或

的每線程陣列尺寸足夠小，以分配足夠的空間用於多個線程塊中的所有線程（每個線程1024浮點數遠遠多）。

正如njuffa提到的那樣，你的內核只使用2個寄存器的原因是你沒有對內核中的數據做任何有用的事情，死編碼全部被編譯器清除。

來源

2012-08-29 01:04:01 harrism

您建議線程可以使用的reg的數量是有限制的（SM_20爲63）。這是從哪裏來的？設備屬性顯示每個BLOCK（regsPerbBock）的reg的數量限制。 – Doug

它來自體系結構，編譯器負責確保沒有大於生成的二進制代碼中使用的限制的寄存器數量。除了性能方面的原因（例如瞭解寄存器溢出的原因），用戶不必擔心此限制，這就是爲什麼不需要將其列在deviceProps結構中的原因。 – harrism

使用許多寄存器可能是需要的，因爲最大化佔用並不是隱藏延遲的唯一方法。隱藏延遲的另一種方式是指令級並行。有時它是達到最佳性能的唯一途徑。查看瓦西里沃爾科夫[幻燈片]（http://www.cs.berkeley.edu/~volkov/volkov10-GTC.pdf），其中的作者在只有8％的入住率時獲得了高峯表現。 –

如前所述，寄存器（和PTX「參數空間」）不能動態索引。爲了做到這一點，編譯器必須發出代碼，如switch...case塊來將動態索引變成立即數。我不確定它會自動執行。你可以使用固定大小的元組結構和switch...case來幫助它。 C/C++元編程很可能是保持這種代碼易於管理的首選武器。

另外，對於CUDA 4.0，請使用命令行開關-Xopencc=-O3以便除映射到寄存器（參見this post）的純標量（例如數據結構）以外的任何內容。對於CUDA> 4.0，您必須禁用調試支持（沒有-G命令行選項 - 優化僅在禁用調試時發生）。

PTX級別允許更多虛擬寄存器比硬件。那些在加載時映射到硬件寄存器。您指定的寄存器限制允許您設置生成的二進制文件使用的硬件資源的上限。它作爲編譯器的一種啓發式方法，可以在編譯到PTX時決定何時溢出（見下文）寄存器，因此可以滿足某些併發需求（請參閱CUDA文檔中的「啓動邊界」，「佔用」和「併發內核執行」 - 你也可以享受this most interesting presentation）。

對於Fermi GPU，最多有64個硬件寄存器。第64個（或最後一個 - 當使用小於硬件的最大值時）被ABI用作堆棧指針，因此用於「寄存器溢出」（這意味着通過臨時將其值存儲在堆棧中釋放寄存器，並在更多寄存器需要比可用），所以它是不可接觸的。

來源

2012-08-29 20:28:04 Dude

有關-Xopencc = -O3的鏈接消失了，我無法在CUDA上下文中找到任何引用。你能否指點我一些資源或解釋近期cuda（7.0/7.5）的行爲是否相似？ – XapaJIaMnu

強制CUDA使用寄存器作爲變量

回答

相關問題