CPU 超过 GPU

Posted 2023-03-23

【中文标题】CPU 超过 GPU

【英文标题】：CPU overtakes the GPU

【发布时间】：2016-02-19 04:30:33

【问题描述】：

我对理解SM中的CUDA线程处理有些疑惑。从我一直在阅读的内容中推断出以下命题： 我的 GPU 是：GTX650Ti。

块中的线程计数必须始终是 Warp 大小的倍数。因此，每个 SM 可以处理 32 个线程的块 (warpSize)。 我的 SM 可以同时计算的最大线程数为 2048 (maxThreadsPerMultiProcessor)。 由于每个 SM 可以同时计算 2048 个线程，warpSize 为 32，因此可以同时计算 64 个块。 由于我的GPU有4个SM，可以同时执行64X4=256个线程块。 因此，内核启动可能有以下启动参数：>>，每次内核启动会调用8192个线程。

是吗？

因此，如果我的内核中有一个包含 10M 元素的向量要处理，这意味着我必须将其分割为 1221 个作业（内核启动），每个作业包含 8192 个元素？

之所以出现这个问题，是因为我正在比较顺序程序和我的 CUDA 程序之间的时间性能。但我只能看到 CPU 超过了 GPU。我还尝试了最大启动参数，例如 >>。结果非常相似。

那么，我在做什么或配置错误？

这是我正在使用的代码：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <math.h>
#include <time.h>
#include "C:\cdev.h"
#include <thrust/device_vector.h>

using namespace thrust;
using namespace std;

#define N (1024 * 16384)

cdev devices;

__global__ void eucliDist(double *c, double *a, double *b)

    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < N)
        c[i] = sqrt(pow(a[i], 2) + pow(b[i], 2));


int main()

    clock_t start, end;
    double elapsed;
    static double A[N];
    static double B[N];
    for (int i = 0; i < N; i++)
    
        A[i] = double(i);
        B[i] = double(i);
       
    static double C[N];

    // Sequential execution of F(x,y) = sqrt((x^2 + y^2))
    start = clock();
    for (int i = 0; i < N; i++)
        C[i] = sqrt(pow(A[i], 2) + pow(B[i], 2));
    end = clock();
    elapsed = double(end - start) / CLOCKS_PER_SEC;
    cout << "Elapsed time for sequential processing is: " << elapsed << " seconds." << endl;

    // CUDA Initialization
    unsigned int threadNum;
    unsigned int blockNum;
    cudaError_t cudaStatus;
    threadNum = devices.ID[0].maxThreadsPerBlock;
    blockNum = ceil(double(N) / double(threadNum));
    // Parallel execution with Thrust of F(x,y) = sqrt((x^2 + y^2))
    vector<double> vectorA(N);
    vector<double> vectorB(N);
    for (int i = 0; i < N; i++)
    
        vectorA[i] = double(i);
        vectorB[i] = double(i);
    
    vector<double> vectorC(N);
    start = clock();
    device_vector<double> thrustA(N);
    cudaStatus = cudaGetLastError();
    if (cudaStatus != cudaSuccess)
    
        cerr << "Device vector allocation failed: " << cudaGetErrorString(cudaStatus) << " (thrustA)" << endl;
        cin.get();
        return 1;
    
    device_vector<double> thrustB(N);
    cudaStatus = cudaGetLastError();
    if (cudaStatus != cudaSuccess)
    
        cerr << "Device vector allocation failed: " << cudaGetErrorString(cudaStatus) << " (thrustB)" << endl;
        cin.get();
        return 1;
    
    device_vector<double> thrustC(N);
    cudaStatus = cudaGetLastError();
    if (cudaStatus != cudaSuccess)
    
        cerr << "Device vector allocation failed: " << cudaGetErrorString(cudaStatus) << " (thrustC)" << endl;
        cin.get();
        return 1;
    
    thrustA = vectorA;
    cudaStatus = cudaGetLastError();
    if (cudaStatus != cudaSuccess)
    
        cerr << "Host to device copy failed (Thrust): " << cudaGetErrorString(cudaStatus) << " (vectorA -> thrustA)" << endl;
        cin.get();
        return 1;
    
    thrustB = vectorB;
    cudaStatus = cudaGetLastError();
    if (cudaStatus != cudaSuccess)
    
        cerr << "Host to device copy failed (Thrust): " << cudaGetErrorString(cudaStatus) << " (vectorB -> thrustB)" << endl;
        cin.get();
        return 1;
    
    eucliDist <<<blockNum, threadNum>>>(raw_pointer_cast(thrustC.data()), raw_pointer_cast(thrustA.data()), raw_pointer_cast(thrustB.data()));
    cudaStatus = cudaGetLastError();
    if (cudaStatus != cudaSuccess)
    
        cerr << "Kernel launch failed (Thrust): " << cudaGetErrorString(cudaStatus) << " (euclidDist)" << endl;
        cin.get();
        return 1;
    
    thrust::copy(thrustC.begin(), thrustC.end(), vectorC.begin());
    cudaStatus = cudaGetLastError();
    if (cudaStatus != cudaSuccess)
    
        cerr << "Device to host copy failed: " << cudaGetErrorString(cudaStatus) << " (thrustC -> vectorC)" << endl;
        cin.get();
        return 1;
    
    end = clock();
    elapsed = double(end - start) / CLOCKS_PER_SEC;
    cout << "Elapsed time parallel processing is (Thrust): " << elapsed << " seconds." << endl;

    cin.get();
    return 0;

建议将不胜感激。

【问题讨论】：

在分析器中运行你的程序，看看时间花在了哪里。不要挂在硬件细节上，只需启动内核一次，使用尽可能多的块来覆盖整个 10M 元素的网格。从教科书“添加向量”示例开始学习该技术。而且 GPU 并不总是比 CPU 快；这取决于任务。

我会说，在第 1 点的第一句话之后，您列表中的所有内容都不正确...

我可以证明这一点。请参阅 CUDA 编程：Shane Cook 的 GPU 并行计算开发人员指南，第 83 页，GRIDS 部分，第 2 段，第 1 行：“块中的线程数应始终是 Warp 大小的倍数，目前定义为 32。”当我说必须时，我可能夸大了，所以这只是建议不要使用条件语句来处理向量大小的元素。

您在很多方面感到困惑，并有效地要求在这个问题中提供有关许多 CUDA 主题的广泛教程。正如@talonmies 所说，您几乎所有的猜想都是错误的。从哪儿开始？也许你应该研究一个简单的代码，比如 vectorAdd 并观察它可以在单个内核启动中处理一个大向量。您绝对不需要 1221 内核启动来处理 10M 元素向量。

@Vitrion：但为什么它应该自动更快呢？您完全没有提供有关您的代码的任何信息，因此无法说出您的代码可能存在什么性能问题或您可能会采取什么措施来修复它们。

【参考方案1】：

让我们首先更正您在问题中发布的很多内容：

块中的线程数应始终是 Warp 大小的倍数。 每个 SM 可以处理 32 个线程 (warpSize) 的倍数的块，每个块最多可以处理 1024 个线程 (cudaDevAttrMaxThreadsPerBlock)。 我的 SM 可以同时计算的最大线程数为 2048 (cudaDevAttrMaxThreadsPerMultiProcessor)。 一个 SM 上最多可以同时驻留 16 个块。 由于我的 GPU 有 4 个 SM，最多可以有 16 x 4 = 64 个线程块 同时执行。 内核启动参数可以是architecture maximum 之前的任何参数，但受制于here 汇总的资源限制。设备上的最大常驻线程数为 4 x 2048 = 8192 个线程。

因此，如果我的内核中有一个包含 10M 元素的向量要处理，这意味着我必须将其分割为 1221 个作业（内核启动），每个作业包含 8192 个元素？

不，您将在一次内核启动中启动 9766 个 1024 个线程的块。或者启动足够的块来完全占用您的 GPU（最多 64 个，具体取决于资源），并让每个线程处理输入向量的多个元素。

【参考方案2】：

您应该分解每个操作的时间安排；您可能对每个元素做的工作太少，以至于您将大部分时间花在主机和设备之间来回复制内存。

如果计算确实是问题所在，那可能是您尝试执行的操作。 pow(x,2) 不是一种特别有效的求平方数的方法。虽然这在 CPU 上很糟糕，但在 GPU 上尤其糟糕，因为这可能意味着您必须使用特殊功能单元，而且这些功能单元并不多，因此会造成瓶颈，因为这是您计算的大部分。

（除此之外：单精度（倒数）平方根在不同的功能单元中处理，该单元具有更多可用的吞吐量）

更糟糕的是，您使用的是双精度浮点数； GPU 设计用于处理单精度浮点数。虽然它可以执行双精度，但它的吞吐量要低得多。

因此，你应该

用x*x 而不是pow(x,2) 计算平方 使用float 而不是double（如果适合您的应用程序）

【讨论】：

是的，我会试试的。但我想告诉你一件事。我尝试删除 pow 和 sqrt 函数并仅执行简单的求和。结果是一样的：CPU 超过了 GPU。我一直在尝试逐步运行它，并且我意识到从主机复制到设备时会出现问题，反之亦然。你知道为什么会这样吗？我怀疑使用推力::copy() 会降低性能。

@Vitrion：那么计算无关紧要：您的基准测试只是测量内存带宽。主机内存和 CPU 之间的带宽大于主机内存和设备内存之间的带宽。

@Vitrion - 当您将double、sqrt 和pow 全部放在一个位置时，GPU 计算吞吐量会很低。当您不进行或进行微不足道的计算（例如简单的添加）时，您会浪费时间在主机和设备之间发送数据。您需要在 GPU 上进行 大量 高吞吐量 计算，以超过 PCIe 传输开销。就像已经说过的那样，GPU 不仅仅是自动比 CPU 快。

【参考方案3】：

我的应用程序使用thrust::device_vector 在设备中分配内存。这就是试图在我的程序中使用推力的原因。 我终于找到了提高 GPU 性能优于 CPU 的问题和解决方案。这对于决定使用 device_vectors 而不是数组的其他用户很有用。

正如@Hurkyl 对我所说：我正在测量主机和设备之间复制的延迟，反之亦然。所有这些长延迟都是由于使用了以下说明：

thrustA = vectorA 用于从主机复制到设备。此复制操作可能会产生清晰而优雅的结果，但请小心。 thrust::copy 用于从设备复制到主机。该功能与使用带有 std::vector 复制功能的复制非常相似。

这两个操作是我代码中的瓶颈。

考虑变量：

vector<double> A;
device_vector<double> thrustA;

解决方案非常简单。我只是用众所周知的 cudaMemcpy() 函数替换了这两条指令，即

从主机复制到设备：

cudaMemcpy(raw_pointer_cast(thrustA.data()), raw_pointer_cast(A.data()), A.size(), cudaMemcpyHostToDevice);

对于从设备复制到主机：

cudaMemcpy(raw_pointer_cast(A.data()), raw_pointer_cast(thrustA.data()), A.size(), cudaMemcpyDeviceToHost);

感谢所有花时间解决我的问题的人。您的意见非常丰富，让我对 CUDA 有了更深入的了解。