c++ - 为什么一个简单的 CUDA 函数需要这么多本地内存？

Question

我在 CUDA 上写了一个简单的函数。它将图像的大小调整为两倍。对于 1920*1080 的图像，此功能需要 ~20ms 才能完成。我尝试了一些不同的方法来优化该功能。而且我发现可能是本地内存是关键原因。

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

• OpenCV中的Gpu模块
• 纹理绑定(bind)到 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 测试的，本地内存负载为零)？

我的代码是在 Debug模式下编译的。而我的卡是GT650(192 SP/SM, 2 SM)。

Answer 1

您发布的两个函数太简单了，无法使用那么多堆栈，事实上它们根本不应该使用堆栈。溢出过多的最可能原因是您在禁用优化的情况下进行编译(例如，在 Debug模式下)。

作为引用，Robert Crovella 在发布和 Debug模式下编译了您的第一个内核:

调试:

ptxas info : Function properties for _Z18performDoubleImagePfmii 256 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 23 registers, 296 bytes cumulative stack size, 56 bytes cmem[0], 1 textures

发布:

ptxas info : Function properties for _Z18performDoubleImagePfmii 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 9 registers, 56 bytes cmem[0], 1 textures

注意堆栈和寄存器使用的区别。如评论中所述，在测量程序性能时，您应该始终针对最大优化级别进行编译，否则测量将毫无意义。