CUDA和openCV (CPU)矩阵加法性能常数随数值的增加

Question

通过连续增加元素数和测量运行时，比较了使用简单CPU功能的矩阵加法、CUDA和openCV (on )的性能。我已经绘制了下面的数据。请注意，它是每个数据类型的一个地块，其中CUCV_8U是用于无符号字符、CUCV_16U=unsigned short、CUCV32F=float和CUCV64F=double的宏。

我注意到openCV和CUDA的运行时不会增加，直到矩阵大约有2^12元素。在超出“极限”之后，运行时开始发散(注意对数缩放)。现在，我想解释一下这个“限制”。

如果它只是为CUDA，我认为这是由于有数量的可用的CUDA核心，这是1024我的GTX 960。当矩阵中的总元素数超过核心数时，线程不能再并行执行，而是并行执行。

然而，openCV似乎遵循着同样的趋势(持续运行直到大约2^12元素)，所以我想知道这是什么原因。有人有主意吗？

谢谢!

如果您需要有关我使用的代码的信息，请查看我的项目回购https://github.com/Vinc37-git/cuCV或下面的内容。

​

​

// CUDA Kernel

template <typename T>
__global__ void cuCV::kernel::add(DeviceCuMat<T> OUT, const DeviceCuMat<T> A, const DeviceCuMat<T> B) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int  ch = blockIdx.z * blockDim.z + threadIdx.z;

    int index = row * A.getWidth() + col + (A.getWidth()*A.getHeight()) * ch;  // linearisation of index

    if (col < A.getWidth() && row < A.getHeight() && ch < A.getNChannels())
        OUT.getDataPtr()[index] = A.getDataPtr()[index] + B.getDataPtr()[index];
}

Answer 1

我注意到openCV和CUDA的运行时不会增加，直到矩阵大约有2^12元素。

启动内核需要一段时间，因为它需要与操作系统(更具体地说，是图形驱动程序)进行交互，而目标设备通常是PCI设备(需要PCI通信)。I/O操作通常是相当慢的，它们通常有相当大的延迟(对于HPC应用程序)。还有许多其他的开销需要考虑，包括分配设备数据的时间、管理虚拟内存页等。由于系统调用在大多数主流系统上至少需要占用1我们，对于涉及离散GPU的系统大约需要花费我们数十人，因此报告的时间安排是预期的。启动数据传输到基于PCI的GPU的时间通常是数百人(几乎所有平台上至少有10人)。

请注意，GPU是为执行大规模并行计算代码而设计的，而不是低延迟代码。拥有更多的核心并不能使小的计算更快，但更多的操作可以计算得更快。这就是安达尔氏和古斯塔夫森氏法律的原则。

如果要计算小矩阵，则不应使用离散GPU。如果您有很多它们，那么您可以在GPU上高效地这样做，但是使用一个大内核(而不是每个矩阵一个内核)。

如果只是针对CUDA，我认为这是因为我的GTX 960有1024个可用的CUDA核心。

这样的计算不应该是计算绑定的，而应该是memory-bound (否则，您应该明确地优化内核，因为任何GPU都应该为这样的基本操作填充内存)。因此，重要的是目标GPU的内存带宽，而不是CUDA核心的数量。还要注意的是，GPU的内存速度往往比CPU所能使用的要快，但延迟要大得多。

还要注意的是，GTX 960的设计并不是为了高效地计算双精度计算.事实上，在这样的计算中，一个好的主流CPU显然应该优于它。这个GPU是为了加速简单精确的计算(通常用于像游戏这样的3D应用程序)。如果您想要加快双精度计算，那么您需要一个(更昂贵的)基于服务器的GPU。幸运的是，对于内存绑定代码来说，这应该不是什么大问题。

OpenCV似乎遵循着同样的趋势(持续运行直到大约2^12元素)。

OpenCV实现仍然存在开销。例如，创建线程需要一些时间，也需要分配数据(尽管它应该比GPU上快得多)，更不用说缓存也是冷的(因此可能受到缓存错误的约束)。在主流Linux上创建几个线程的时间大约是很少到几十个人(由于更多的核心和更复杂的体系结构，服务器上的线程通常更多)。

总的来说，CPU实现似乎是内存绑定的，因为计算的速度取决于项大小(乍一看是成比例的)。尽管如此，假设结果为一个内核执行的代码(我得到了2.5GIB/s，这对于最近的主流PC来说是相当糟糕的)，那么吞吐量就显得相当缓慢。GPU的计算效率也很低，因为时间与项目大小无关。这意味着计算当然不受内存限制，内核可能会被优化(检查映像的布局，GPU首先进行连续访问，然后应用一些展开，以减少每个线程内核的开销)。

我强烈建议您使用配置文件( CPU )和GPU实现。CPU实现的第一件事(对于大型矩阵)是检查是否使用多个核，以及是否执行SIMD指令。对于CPU和GPU实现，您应该检查内存吞吐量。