为什么 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 个我能想到的不同因素。但是，通过将测量与实际计算隔离开来，而不是数据传输或“完整操作”，您已经在某种程度上简化了问题。

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

内核启动开销 - 在 GPU 上启动内核会带来“大约”固定成本开销，每次内核启动通常在 5 到 50 微秒的范围内。这意味着，如果您调整工作量以使 CPU 可以在不到该时间的时间内完成，GPU 就不可能更快。甚至在该级别之上，还有一个描述该开销模型的线性函数，我相信如果您愿意，您可以计算出该模型，以在存在固定成本开销的情况下比较 CPU 与 GPU 性能。在比较小的测试用例时，这是一个需要考虑的重要因素，但我的猜测是，因为大多数测试用例的时间都远高于 50 微秒，我们可以放心地“忽略”这个因素，作为一个近似值。
实际CPU与实际GPU的实际performance/capability。这通常很难建模，取决于您使用的特定硬件，并且您没有提供该信息。然而，无论如何我们都可以根据您提供的数据进行一些观察和一些推测，在下一节中对此进行扩展。

考虑到 N=5000 和 N=1000000，您的两个案例涉及 N 描述的总工作量。构建一个小图表：

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

所以我们看到，在 CPU 的情况下，当工作量增加 200 倍时，执行时间增加了 ~352 倍，而在 GPU 的情况下，执行时间增加了约 4.5 倍。我们需要解释这两个“non-linearities”，以便合理猜测发生了什么。

缓存的影响 - 因为您运行您的测试用例 100 次，所以缓存可能会产生影响。在 CPU 的情况下，这是我对您为什么看不到线性关系的唯一猜测。我猜想在非常小的情况下，你在一些 CPU “内部”缓存中，“可见”有 40KB 的数据。转到更大的大小，您可以看到 8MB 的数据，尽管这可能适合 CPU 上的“外部”缓存，但也可能不适合，即使适合，外部缓存也可能会产生整体性能比内部缓存慢。我想这就是 CPU 随着数据变大而变得更糟的原因。您的 CPU 受到较大数据集 non-linearly 的 负面影响。在 GPU 的情况下，外部缓存最多为 6MB（除非你在 Ampere GPU 上运行），因此你的较大数据集不能完全放入外部缓存中。
机器饱和的影响 - CPU 和 GPU 都可以完全“加载”或部分加载，具体取决于工作负载。在 CPU 的情况下，我猜您没有使用任何 multi-threading，因此您的 CPU 代码仅限于单核。（而且，您的 CPU 几乎肯定有多个可用内核。）您的单线程代码将大致“饱和”，即保持该单核“忙碌”。然而，GPU 有很多内核，我猜你的小测试用例（将达到 5000 个线程）只会使你的 GPU 部分饱和。我的意思是，在较小的情况下，一些 GPU 线程处理资源将处于空闲状态（除非您碰巧在最小的 GPU 上运行）。 5000 个线程只够让 2 个 GPU SM 保持忙碌，所以如果你的 GPU 有超过 2 个 SM，它的一些资源在较小的测试用例中是空闲的，而你的 million-thread 较大的测试用例足以饱和，即在任何当前的 CUDA GPU 上保持所有线程处理资源繁忙。这样做的效果是，虽然 CPU 根本无法从更大的测试用例中受益（您应该考虑使用 multi-threading），但您的 GPU 可能会受益。较大的测试用例允许您的 GPU 在与较小的测试用例相同的时间内完成更多的工作。因此，GPU 从更大的工作负载中以积极的方式受益 non-linearly。

当给定足够大的工作负载时，GPU 也能够更好地减轻外部缓存缺失的影响。这称为 GPU 在存在“大”并行工作负载时的 latency-hiding 效应，而 CPU 没有（或没有那么多）相应的机制。因此，根据您的确切 CPU 和 GPU，这可能是一个额外的因素。我不打算在这里给出关于 latency-hiding 的完整教程，但这个概念是基于部分的关于上面的第 2 项，因此您可以从中收集一般 idea/benefit。

为什么 cpu 对于小输入比 gpu 快？

Why is the cpu faster than the gpu for small inputs?

cpu

cuda

gpu