为什么简单的CUDA功能需要如此多的本地内存？

Question

我在CUDA上写了一个简单的函数.它将图像调整为双倍缩放.对于1920*1080的图像,此功能需要~20ms才能完成.我尝试了一些不同的方法来优化该功能.而我发现可能是本地记忆是关键原因.

我尝试了三种不同的方法来获取图像.

• OpenCV中的Gpu模块
• 纹理绑定到OpenCV中的GpuMat
• 直接从全局内存中获取GpuMat

他们都不能给我带来一点改善.

然后我使用nvvp找出原因.在上述所有三个条件下,本地内存开销约为95%.

所以我转向我的代码,了解nvcc如何使用内存.然后我发现一个简单的函数就像这样:

__global__ void performDoubleImage(float* outData, size_t step, const int cols, const int rows)
{
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    if (x >= cols)
        return;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    if (y >= rows)
        return;
    ((float*)((size_t)outData+y*step))[x] = tex2D(texRef, x/2, y/2);
}

需要80字节的堆栈帧(它们在本地内存中).

而另一个功能是这样的:

__global__ void performFinalDoubleImage(const PtrStepSz<float> in, PtrStepSz<float> out)
{
    out(out.rows-1, out.cols-1) = out(in.rows-1, in.cols-1);
}

还需要88字节的堆栈帧.

问题是,为什么我的函数在这个简单的任务中使用了如此多的本地内存和寄存器？为什么OpenCV中的函数可以通过不使用本地内存来执行相同的功能(这是由nvvp测试,本地内存负载是ZERO)？

我的代码是在调试模式下编译的.我的卡是GT650(192 SP/SM,2 SM).

Answer 1

你已经发布的两个功能实在是太简单被使用的堆栈,其实他们不应该使用堆栈在所有.出现这么多溢出的最可能原因是您正在编译时禁用了优化(例如,在调试模式下).

作为参考,Robert Crovella在发布和调试模式下编译了您的第一个内核:

调试:

ptxas info:_Z18performDoubleImagePfmii的函数属性256字节堆栈帧,0字节溢出存储,0字节溢出加载ptxas info:使用23个寄存器,296字节累积堆栈大小,56字节cmem [0],1个纹理

发布:

ptxas info:_Z18performDoubleImagePfmii的函数属性0字节堆栈帧,0字节溢出存储,0字节溢出加载ptxas info:使用9个寄存器,56字节cmem [0],1个纹理

注意堆栈和寄存器使用的差异.如评论中所述,在测量程序的性能时,应始终编译以获得最大优化级别,否则测量将毫无意义.