为什么小输入的 cpu 比 gpu 快？

Question

我曾经历过，对于小输入大小，CPU 的执行速度比 GPU 快。为什么是这样？准备，数据传输还是什么？

例如对于内核和 CPU 功能（CUDA 代码）：

__global__ void squareGPU(float* d_in, float* d_out, unsigned int N) {
    unsigned int lid = threadIdx.x;
    unsigned int gid = blockIdx.x*blockDim.x+lid;
    if(gid < N) {
        d_out[gid] = d_in[gid]*d_in[gid]; 
    }
}

void squareCPU(float* d_in, float* d_out, unsigned int N) {
    for(unsigned int i = 0; i < N; i++) {
        d_out[i] = d_in[i]*d_in[i]; 
    }
}

在 5000 个 32 位浮点数的数组上运行这些函数 100 次，我使用一个小测试程序得到以下结果

Size of array:
5000
Block size:
256

You chose N=5000 and block size: 256

Total time for GPU: 403 microseconds (0.40ms)
Total time for CPU: 137 microseconds (0.14ms)

将数组的大小增加到 1000000，我得到：

Size of array:
1000000
Block size:
256

You chose N=1000000 and block size: 256

Total time for GPU: 1777 microseconds (1.78ms)
Total time for CPU: 48339 microseconds (48.34ms)

我不包括用于在主机和设备之间传输数据的时间（反之亦然），事实上，这是我的测试程序的相关部分：

gettimeofday(&t_start, NULL);

for(int i = 0; i < 100; i++) {
    squareGPU<<< num_blocks, block_size>>>(d_in, d_out, N);
} cudaDeviceSynchronize();

gettimeofday(&t_end, NULL);

选择块大小后，我计算相对于数组大小的块数： unsigned int num_blocks = ((array_size + (block_size-1)) / block_size);

Answer 1

回答 CPU 与 GPU 性能比较的一般问题是相当复杂的，并且通常涉及我能想到的至少 3 或 4 个不同因素的考虑。但是，您通过将测量与实际计算（而不是数据传输或“完整操作”）隔离开来，在某种程度上简化了问题。

在这种情况下，可能至少需要考虑两件事：

1. 内核启动开销 - 在 GPU 上启动内核会带来“大约”固定成本开销，每次内核启动通常在 5 到 50 微秒的范围内。这意味着，如果您调整工作量的大小，使您的 CPU 可以在少于该时间量的时间内完成工作，那么 GPU 就不可能更快。即使在该级别之上，也有一个线性函数描述了开销模型，我相信如果您愿意，您可以使用它来比较 CPU 与 GPU 性能在存在固定成本开销的情况下。在比较小型测试用例时，这是一个需要考虑的重要因素，但我的猜测是，因为您的大多数测试用例时间都远高于 50 微秒，我们可以安全地“忽略”这个因素，作为近似值。

2. 实际 CPU 与实际 GPU 的实际性能/能力。这通常很难建模，取决于您使用的特定硬件，并且您尚未提供该信息。但是，无论如何我们都可以根据您提供的数据进行一些观察和一些猜测，并在下一节对此进行扩展。

你的两个案例涉及的工作总量由N、 、N=5000和描述N=1000000。建立一个小图表：

      N  |  CPU time    |  GPU time
   5000  |    137       |  403
1000000  |  48339       | 1777

所以我们看到，在 CPU 的情况下，当工作量增加 200 倍时，执行时间增加了约 352 倍，而在 GPU 情况下，执行时间增加了约 4.5 倍。我们需要解释这两种“非线性”，以便对正在发生的事情有一个合理的猜测。

1. 缓存的影响 - 因为您运行了 100 次测试用例，所以缓存可能会产生影响。在 CPU 情况下，这是我对为什么没有看到线性关系的唯一猜测。我猜想，在非常小的尺寸下，您处于某个 CPU“内部”缓存中，“查看”了 40KB 的数据。使用更大的尺寸，您可以查看 8MB 的数据，尽管这可能适合您 CPU 上的“外部”缓存，但也可能不适合，即使适合，外部缓存的整体性能也可能比内部缓存。我猜这就是 CPU 看起来随着数据变大而变得更糟的原因。您的 CPU 受到非线性影响负面影响方式，来自更大的数据集。在 GPU 情况下，外部缓存最多为 6MB（除非您在 Ampere GPU 上运行），因此您的较大数据集无法完全放入外部缓存中。

2. 机器饱和的影响 - CPU 和 GPU 都可以完全“加载”或部分加载，具体取决于工作负载。在 CPU 情况下，我猜您没有使用任何多线程，因此您的 CPU 代码仅限于单核。（而且，您的 CPU 几乎肯定有多个可用内核。）您的单线程代码将近似“饱和”，即保持该单核“忙碌”。然而，GPU 有很多核心，我猜你的较小的测试用例（可以达到 5000 个线程）只会部分饱和你的 GPU。我的意思是一些 GPU 线程处理资源在较小的情况下将处于空闲状态（除非您碰巧在最小的 GPU 上运行）。5000 个线程只够让 2 个 GPU SM 忙碌，所以如果你的 GPU 有超过 2 个 SM，它的一些资源在较小的测试用例中处于空闲状态，而您的百万线程较大的测试用例足以饱和，即在任何当前的 CUDA GPU 上保持所有线程处理资源忙碌。这样做的结果是，虽然 CPU 根本无法从更大的测试用例中受益（您应该考虑使用多线程），但您的 GPU 可能会受益。更大的测试用例可以让你的 GPU 做更多的工作在与较小的测试用例相同的时间内。因此，GPU 从更大的工作负载中以积极的方式非线性地受益。

当 GPU 被赋予足够大的工作负载时，它也能够更好地减轻外部缓存中丢失的影响。这称为 GPU 在存在“大型”并行工作负载时的延迟隐藏效应，而 CPU 没有（或没有那么多）相应的机制。因此，根据您的确切 CPU 和 GPU，这可能是一个额外的因素。我不打算在此处提供有关延迟隐藏的完整教程，但该概念部分基于上面的第 2 项，因此您可以从中收集总体思路/好处。