2

我在 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)。

4

1 回答 1

6

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

作为参考，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 个纹理

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

于 2014-12-12T16:33:50.930 回答