我遵循了 CUDA 教程，但我的 GPU 计算时间比我的 CPU 时间长得多？

Question

我按照this page 上的教程进行操作，但结果很糟糕。所用时间如下：

CPU: 569
GPU: 11160

这是我的代码。出了什么问题？我不明白为什么这段代码这么慢？

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <chrono>
#include <iostream>
#include <math.h>
#include <stdio.h>

__global__ void addCUDA(int n, float* x, float* y)
{
    int index = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;
    for (int i = index; i < n; i += stride)
        y[i] = x[i] + y[i];
}

void add(int n, float* x, float* y)
{
    for (int i = 0; i < n; i++)
        y[i] = x[i] + y[i];
}

int main()
{
    int N = 1 << 20;
    float* x = new float[N];
    float* y = new float[N];
    for (int i = 0; i < N; i++) {
        x[i] = 1.0f;
        y[i] = 2.0f;
    }
    auto t1 = std::chrono::high_resolution_clock::now();
    add(N, x, y);
    auto t2 = std::chrono::high_resolution_clock::now();
    auto duration = std::chrono::duration_cast<std::chrono::microseconds>(t2 - t1).count();
    float maxError = 0.0f;
    for (int i = 0; i < N; i++)
        maxError = fmax(maxError, fabs(y[i] - 3.0f));
    std::cout << "Max error: " << maxError << std::endl;
    std::cout << duration << std::endl;
    delete[] x;
    delete[] y;

    float* u,
    float* v;
    cudaMallocManaged(&u, N * sizeof(float));
    cudaMallocManaged(&v, N * sizeof(float));
    for (int i = 0; i < N; i++) {
        u[i] = 1.0f;
        v[i] = 2.0f;
    }
    int blockSize = 256;
    int numBlocks = (N + blockSize - 1) / blockSize;
    int device = -1;
    cudaGetDevice(&device);
    cudaMemPrefetchAsync(u, N * sizeof(float), device, NULL);
    cudaMemPrefetchAsync(v, N * sizeof(float), device, NULL);

    auto t3 = std::chrono::high_resolution_clock::now();
    addCUDA<<<numBlocks, blockSize>>> (N, u, v);
    cudaDeviceSynchronize();
    auto t4 = std::chrono::high_resolution_clock::now();
    duration = std::chrono::duration_cast<std::chrono::microseconds>(t4 - t3).count();
    maxError = 0.0f;
    for (int i = 0; i < N; i++)
        maxError = fmax(maxError, fabs(v[i] - 3.0f));
    std::cout << "Max error: " << maxError << std::endl;
    std::cout << duration << std::endl;
    cudaFree(u);
    cudaFree(v);

    return 0;
}

Answer 1

对于如此微不足道的操作（每个元素上的+），将缓冲区从主机发送到 gpu 并将缓冲区从 gpu 检索到主机需要比执行实际计算更多的时间。 即使 API 非常舒适地使缓冲区访问看起来很容易并且几乎是神奇的，但数据必须通过 PCI-express 总线传输...... 这里的传输是异步的，但是计算必须等待它完成才能真正开始；仅当您同时有其他事情要做时（例如，将复杂计算的各个阶段组织为管道），异步传输才有意义。

如果您尝试解决另一个需要更多计算的问题，缓冲区传输将被摊销。

此外，1<<20floats 的两个数组只需要 8MB，并且可以放入现代 CPU 的高速缓存中。 然后，在这两个数组初始化之后，它们可能已经在高速缓存中热，并且很容易被 CPU 计算访问。 因为计算是一个完全规则的循环，一个体面的优化编译器会使用 SIMD 指令，CPU 不会错误预测分支，并且会完美地预取各种缓存级别中的数据；所有这些都大大提高了这种计算的 CPU 效率。

使用 GPU 超越现代 CPU 并非易事。 这实际上取决于问题的大小和复杂性（当然也取决于这两个硬件的具体属性）。

---

编辑

正如 cmets 中所讨论的，引用文章中使用的计时方法与问题中显示的方法非常不同。

在文章中，nvprof 使用 GPU 中的内部计数器来测量主动计算 addCUDA()（文章中的add()）函数所花费的时间，而不考虑获取两个源缓冲区所需的时间从主机并将生成的缓冲区发送回主机。 当然，速度很快！因为在许多现代硬件（CPU 或 GPU）上，大部分时间都花在访问/传输数据而不是计算上。如果我们测量 CPU 中花费的时间仅，忽略从/向缓存/内存获取/写入数据所花费的时间，也不会很长！ （注意文中的CPU代码甚至没有开启优化编译，这样的计时有什么意义吗？）

在问题中显示的代码中，计时方法完全不同，但在我看来更相关。 对std::chrono::high_resolution_clock::now() 的两次调用实际上考虑了完成所有工作所花费的时间：发送两个源缓冲区，对它们进行计算并获取结果缓冲区。 毕竟，这是唯一重要的持续时间！ 这样，将这个持续时间与我们在对 CPU 计时时（使用类似方法）获得的持续时间进行比较是公平的。

使用cudaMemPrefetchAsync() 的事实可能会产生误导，因为我们可能认为源缓冲区的传输被排除在时间之外：事实并非如此，这就是为什么我们发现与文章相比结果令人失望的原因。 我们在这两个调用之后立即启动计时器以测量计算花费的时间，但计算必须等待这些传输完成才能真正开始（我什至会在 before这两个电话）。 此外，在停止计时器之前对cudaDeviceSynchronize() 的调用会等待结果缓冲区的传输完成，以便实际将结果提供给主机。 如果我们在启动计时器之前使用cudaDeviceSynchronize()，我们可以从计时中排除两次初始传输，但是这样的计时有什么意义呢？

总之，我认为您在问题中使用的计时方法比文章中介绍的方法要好得多，因为您可以真正比较从一种技术与另一种技术获得（或没有！）的好处。

有关信息，在我的计算机上，打开完全优化后，您的代码会给出以下结果：

CPU: 809   Intel(R) Xeon(R) CPU E5-2697 v2 @ 2.70GHz]
GPU: 1160  NVIDIA Corporation GM200 [GeForce GTX TITAN X] (rev a1)

CPU: 157   Intel(R) Core(TM) i7-10875H CPU @ 2.30GHz
GPU: 1158  NVIDIA Corporation TU104GLM [Quadro RTX 4000 Mobile / Max-Q] (rev a1)