cuda註冊壓力

我有一個內核做線性最小二乘擬合。原來線程使用的寄存器太多，因此佔用率很低。這裏是內核，cuda註冊壓力

__global__ 
void strainAxialKernel(
    float* d_dis, 
    float* d_str 
){ 
    int i = threadIdx.x; 
    float a = 0; 
    float c = 0; 
    float e = 0; 
    float f = 0; 
    int shift = (int)((float)(i*NEIGHBOURS)/(float)WINDOW_PER_LINE); 
    int j; 
    __shared__ float dis[WINDOW_PER_LINE]; 
    __shared__ float str[WINDOW_PER_LINE]; 

    // fetch data from global memory 
    dis[i] = d_dis[blockIdx.x*WINDOW_PER_LINE+i]; 
    __syncthreads(); 

    // least square fit 
    for (j=-shift; j<NEIGHBOURS-shift; j++)          
    {                   
     a += j;                 
     c += j*j;                
     e += dis[i+j];               
     f += (float(j))*dis[i+j];            
    }                  
    str[i] = AMP*(a*e-NEIGHBOURS*f)/(a*a-NEIGHBOURS*c)/(float)BLOCK_SPACING;  

    // compensate attenuation 
    if (COMPEN_EXP>0 && COMPEN_BASE>0)           
    {                   
     str[i]                 
     = (float)(str[i]*pow((float)i/(float)COMPEN_BASE+1.0f,COMPEN_EXP));  
    } 

    // write back to global memory 
    if (!SIGN_PRESERVE && str[i]<0)            
    {                   
     d_str[blockIdx.x*WINDOW_PER_LINE+i] = -str[i];       
    }                   
    else                   
    {                   
     d_str[blockIdx.x*WINDOW_PER_LINE+i] = str[i];       
    } 
}

我有32x404塊，每個塊有96個線程。在GTS 250上，SM應該能夠處理8個塊。然而，可視化剖析器顯示每個線程有11個寄存器，因此佔用率爲0.625（每個SM 5個塊）。順便說一句，每個塊使用的共享內存是792 B，所以寄存器是問題。表演並非世界末日。我只是好奇，如果有任何我可以解決這個問題。謝謝。

來源

2010-11-16 small_potato

怎麼樣的網格配置？ – fabrizioM 2010-11-17 00:14:04

我忘記了，現在修復了 – 2010-11-17 00:28:10

快速但有限的寄存器/共享內存與緩慢但大型的全局內存之間總是存在權衡。沒有辦法「折中」這種交換。如果通過使用全局內存來使用減少寄存器的使用量，則應該獲得更高的佔用率，但訪問速度會更慢。

也就是說，這裏有一些想法，以使用更少的寄存器：

可以轉移預先計算和存儲在常量內存？然後每個線程只需要查找shift [i]。
做一個和c必須漂浮？
或者，可以將a和c從循環中移除並計算一次？並因此徹底刪除？

被計算爲一個簡單的等差數列，所以減少它......（像這樣）

a = ((NEIGHBORS-shift) - (-shift) + 1) * ((NEIGHBORS-shift) + (-shift))/2

或

a = (NEIGHBORS + 1) * ((NEIGHBORS - 2*shift))/2

所以代替，這樣做以下（您可以進一步減少這些表達式）：

str[i] = AMP*((NEIGHBORS + 1) * ((NEIGHBORS - 2*shift))/2*e-NEIGHBOURS*f) 
str[i] /= ((NEIGHBORS + 1) * ((NEIGHBORS - 2*shift))/2*(NEIGHBORS + 1) * ((NEIGHBORS - 2*shift))/2-NEIGHBOURS*c) 
str[i] /= (float)BLOCK_SPACING;

來源

2010-11-18 22:10:39 jmilloy

佔用率不是問題。

GTS 250（計算能力1.1）中的SM可能能夠同時在其寄存器中保存8個塊（8×96個線程），但它只有8個執行單元，這意味着只有8個（或者在你的情況下，5x96）線程將在任何特定時刻推進。嘗試將更多塊壓縮到超載SM上幾乎沒有什麼價值。

實際上，您可以嘗試使用-maxrregcount選項來增加寄存器的數量，這可能會對性能產生積極影響。

來源

2010-11-18 23:03:10 user434507

增加佔用率可以讓每個SM發出更多的內存請求。性能幾乎總是存儲在GPU上。在增加SM上的塊數方面有很多價值！他們幾乎總是坐在無聊之中，因爲他們正在等待數據來自記憶。 – mch 2010-11-21 15:51:04

可以使用啓動界限指示編譯器爲每個多處理器的最大線程數和最小塊數生成寄存器映射。這可以減少寄存器數量，以便達到所需的佔用率。

對於您的情況，Nvidia的佔用率計算器顯示了63％的理論高峯佔用率，這似乎是您的成就。這是由於你的註冊計數，正如你提到的，但這也是由於每塊的線程數。將每個塊的線程數量增加到128個，並將寄存器數量減少到10個，可獲得100％的理論峯值佔用率。

要控制新開工界限爲你的內核：

__global__ void 
__launch_bounds__(128, 6) 
MyKernel(...) 
{ 
    ... 
}

然後只需用128個線程塊大小啓動並享受您的入住。編譯器應該生成內核，使其使用10個或更少的寄存器。

來源

2012-04-23 17:26:17

cuda註冊壓力

回答

相關問題