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注册压力

回答

相关问题