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

【问题标题】：Why a simple CUDA function needs so much local memory?为什么一个简单的 CUDA 函数需要这么多本地内存？
【发布时间】：2015-02-11 07:49:06
【问题描述】：

我在 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测试，本地内存负载为零）？

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

【问题讨论】：

那些函数不可能使用那么多堆栈，它们甚至根本不应该使用堆栈。这是什么GPU？您确定您正确解释了 nvvp 输出吗？您是在调试模式下编译还是在禁用优化（即使那样......）或命令行中的任何特殊情况下进行编译？
我的卡是GT650，有192 SP/SM和2 SM。我用调试编译了我的代码。可能是调试模式是什么原因？我不能给你 nvcc 输出，因为我现在在家，明天我会在我的问题中添加额外的信息，谢谢！
你如何确定堆栈帧？我将你的第一个内核放入一个随机 cuda 程序中，当我使用 -Xptxas -v 编译时，我得到：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
@cs512：在调试模式下编译可能是堆栈帧使用的来源。为了使变量始终可跟踪，它们可能需要存储在本地内存中。
当我编译相同的代码并添加-G，我得到：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，所以，是的，调试（-G）似乎有很大的不同。由于这个问题的最初动机似乎是关于性能，因此值得说明的是，如果使用 -G 编译，分析设备代码的性能绝不是一个好主意

标签： c++ c opencv cuda gpu-local-memory

【解决方案1】：

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

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

调试：

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

发布：

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

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

【讨论】：