cuda寄存器压力

Question

我有一个内核进行线性最小二乘拟合。事实证明，线程使用了太多寄存器，因此占用率很低。这是内核，

__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 个块）。 BTW，每个块使用的共享内存是792 B，所以寄存器是问题所在。 演出并不是世界末日。我只是好奇是否有办法可以解决这个问题。谢谢。

Answer 1

您可以使用启动边界来指示编译器为每个多处理器的最大线程数和最小块数生成寄存器映射。这可以减少寄存器数量，从而达到所需的占用率。

对于您的情况，Nvidia 的占用率计算器显示理论峰值占用率为 63%，这似乎就是您所实现的目标。正如您提到的，这是由于您的寄存器计数造成的，但它也是由于每个块的线程数造成的。将每个块的线程数增加到 128，并将寄存器计数减少到 10，可获得 100% 的理论峰值占用率。

要控制内核的启动范围：

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

然后以 128 个线程的块大小启动并享受您的占用。编译器应该生成您的内核，使其使用 10 个或更少的寄存器。

Answer 2

快速但有限的寄存器/共享内存与缓慢但大的全局内存之间总是存在权衡。没有办法“绕过”这种权衡。如果您通过使用全局内存来减少寄存器的使用，您应该获得更高的占用率，但内存访问速度会更慢。

也就是说，这里有一些使用更少寄存器的想法：

1. 可以预先计算移位并将其存储在常量内存中吗？那么每个线程只需要查找shift[i]即可。
2. a 和 c 必须是浮点数吗？
3. 或者，可以将 a 和 c 从循环中删除并计算一次吗？从而完全删除？

a 被计算为一个简单的算术序列，因此减少它...（类似这样）

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;

Answer 3

入住不是问题。

GTS 250（计算能力1.1）中的SM可能能够在其寄存器中同时保存8个块（8x96线程），但它只有8个执行单元，这意味着8x96中只有8个（或者，在您的情况下，5x96）线程将在任何给定的时间前进。试图将更多的块挤到超载的 SM 上几乎没有什么价值。

事实上，您可以尝试使用 -maxrregcount 选项来增加寄存器的数量，这可能会对性能产生积极的影响。