減少CUDA內核中使用的寄存器的數量

Question

我有一個使用17個寄存器的內核，將其減少到16會帶給我100％的佔用率。我的問題是：是否有方法可用於減少使用的數量或寄存器，不包括以不同方式完全重寫我的算法。我總是認爲編譯器比我聰明得多，所以爲了清晰起見，我經常使用額外的變量。這個想法我錯了嗎？

請注意：我不知道有關--max_registers（或任何語法）標誌，但使用的本地內存會比降低了25％的入住更不利（我應該測試這個）

Answer 1

真的很難說，在我看來，nvcc編譯器不是很聰明。
你可以嘗試一些明顯的事情，例如使用short而不是int，通過引用傳遞和使用變量（例如&變量），展開循環，使用模板（如在C++中）。如果你有分裂，先驗功能，順序應用，嘗試使它們成爲一個循環。儘量擺脫條件，可能用冗餘計算替代它們。

如果你發佈了一些代碼，也許你會得到具體的答案。

Answer 2

入住率可能有點誤導，100％的入住率不應該是您的主要目標。如果可以完全合併訪問全局內存，那麼在高端GPU上佔用50％就足以隱藏全局內存的延遲（對於浮點數，甚至更低）。查看去年GTC的Advanced CUDA C演示文稿，瞭解更多關於此主題的信息。

對於你的情況，你應該測量有無maxrregcount設置爲16的性能。假設你沒有隨機訪問本地數組，那麼本地內存的延遲應該隱藏起來，因爲它有足夠的線程。導致非合併訪問）。

要回答您關於減少寄存器的具體問題，請發佈代碼以獲取更詳細的答案！瞭解編譯器如何在一般情況下工作可能會有幫助，但請記住，nvcc是一個具有大參數空間的優化編譯器，因此最大限度地減少寄存器數量必須與總體性能保持平衡。

Answer 3

降低寄存器使用時的指令數增加有一個簡單的解釋。編譯器可以使用寄存器來存儲通過代碼多次使用的一些操作的結果，以避免重新計算這些值，當被迫使用較少的寄存器時，編譯器決定重新計算將存儲在寄存器中的那些值除此以外。

Answer 4

這通常不是一種好方法來最小化記錄壓力。編譯器在優化整體計劃的內核性能方面做得很好，並且考慮了很多因素，包括註冊表。

它是如何工作的時候降低寄存器造成速度較慢

最有可能的編譯器必須足夠的寄存器中的數據溢出到「本地」的內存，這是基本相同的全局內存，因而非常緩慢

爲了優化目的，我會建議在必要時使用像const，volatile等關鍵字來幫助編譯器優化階段。

無論如何，這不是像寄存器這些微小的問題，經常使CUDA內核運行緩慢。我建議優化全局內存，訪問模式，儘可能在紋理內存中緩存，通過PCIe進行交易。

Answer 5

利用共享內存作爲高速緩存可能會導致更少的註冊使用和防止溢出的寄存器對本地內存...

想想內核計算一些數值與這些計算值是由所有線程的使用，

__global__ void kernel(...) { 
    int idx = threadIdx.x + blockDim.x * blockIdx.x; 
    int id0 = blockDim.x * blockIdx.x; 

    int reg = id0 * ...; 
    int reg0 = reg * a/x + y; 


    ... 

    int val = reg + reg0 + 2 * idx; 

    output[idx] = val > 10; 
} 

所以，與其保持REG和REG0寄存器和使他們possibily溢出到本地內存（全局內存），我們可能會使用共享內存。

__global__ void kernel(...) { 
    __shared__ int cache[10]; 

    int idx = threadIdx.x + blockDim.x * blockIdx.x; 

    if (threadIdx.x == 0) { 
     int id0 = blockDim.x * blockIdx.x; 

     cache[0] = id0 * ...; 
     cache[1] = cache[0] * a/x + y; 
    } 
    __syncthreads(); 


    ... 

    int val = cache[0] + cache[1] + 2 * idx; 

    output[idx] = val > 10; 
} 

看看這個paper瞭解更多信息..