为什么我在 CUDA 中的性能这么差？

Question

我有 Linux 下的 RTX 3050，原生 NVidia 驱动程序。同样的简单代码在 GPU 2.1GHz 上的执行速度比在 CPU 2.4GHz 上慢 10 倍。可能有什么问题：驱动程序、编译器标志，或者这是正常性能？

代码是：

__global__ void testcuda() {
    uint counter = 0;
    do {
        if (++counter % 1000000000 == 0) printf("%u\n", counter);
    } while (counter != 0);
}

testcuda<<<1, 1>>>();

CPU 上为 6 秒，GPU 上为 47 秒

Answer 1

GPU是一个并行机，一个大规模并行机。
如果您查看 GPU 芯片，您会发现它有 2560 个执行核心，而您的 CPU 只有（比如）8 个。

两个芯片占据的表面积大致相同，这意味着 CPU 核心大约是 2560/8 = 320 倍。

所有这些额外的区域都用于减少 CPU 的延迟，通过使用 流水线技术、分支预测、堆栈引擎、乱序执行、超标量执行、推测性缓存预取 ，还有更多。

拥有少量晶体管的 GPU 核心不具备这些功能。相反，GPU 专注于使用类似于 SIMD 的方案并行执行大量语句，NVidia 称之为：SIMT。

这意味着当所有线程对不同数据执行相同指令时，GPU 的工作速度最快。

在您的示例中，您仅运行一个线程，这意味着 2559 个线程处于空闲状态，因此如果您能够完全并行化代码，那么您的运行速度会比原本可以执行的速度慢 2000 倍以上。

并行化代码确实很重要，但如果您能够的话，您将看到 GPU 的好处。

您能实现多大的加速取决于您可以从代码中提取的并行工作量，这受阿姆达尔定律约束。

GPU 线程的正常延迟为 7 或 8 个 GPU 周期，CPU 通常具有 1/2 个周期的延迟（由于流水线、指令重新排序和超标量执行，延迟小于 1）。
考虑到 CPU 有 8 个核心，您必须提取至少 8 x 8 = 64x *) 并行度，然后才能预期 GPU 的性能优于 CPU。

*）这只是一个指示，可能会因应用程序、数据访问模式等而有很大差异。

Answer 2

您仅使用 1 个 cuda 线程，该线程仅使用 1 个扭曲，该扭曲仅属于 GPU 中的 1 个 SM 单元。您的 GPU 有 64 个活动扭曲，甚至更多活动线程的飞行能力。

当您在 CPU 上执行此操作时，它仅损失了 8 倍 - 16 倍的单核性能。但在 GPU 上，它缺少 2048 倍的性能。由于 GPU 管道具有更差的分支预测、更差的堆栈管理、更差的频率和许多其他规格，因此将单个 CUDA 线程性能与单个 CPU 线程性能进行比较是没有用的。单个 GPU 线程仅占 GPU 计算能力的 0.1%。但在 CPU 上，通过一些编译器优化（例如向量化），它可以接近 10%。但 CUDA 编译器不会自动矢量化这些循环以将它们映射到多个管道。 CUDA 编译器只允许您编写单线程代码，并期望您使用数千个线程负责同步、缩减和其他并行原语。因此，您必须使用足够数量的线程来保持 GPU 管道繁忙。

即使您使用 1000 个管道启动了 1000 个线程，所花费的时间也不会少于此时间，因为所有线程都运行相同的内核。您需要显式地将工作分配给线程。大约 1000000 个线程上每个线程 1000 次迭代。也许 1000000 个线程太多了。但 GPU 的线程数量应是管道数量的 16-24 倍才能保持完全占用。

如果 CUDA 具有某种针对简单循环的自动矢量化功能，就像 GCC 那样，它就可以工作。但即便如此，你也没有给它一个适当的循环。由于 printf，该循环中存在依赖链。即使 GCC 也无法自动矢量化 printf 函数。

将 RTX3050 视为堆叠在一起的 Pentium-2000（甚至没有 SSE）集群或 20 年前的超级计算机。如果没有正确的代码，则只能使用超级计算机的单个处理器。