OpenCL 或 CUDA 调用的开销？答案

【问题标题】：The overhead of a OpenCL or CUDA call?OpenCL 或 CUDA 调用的开销？
【发布时间】：2012-02-17 14:51:45
【问题描述】：

我正在编写一个执行大量BLAS gemv 操作的函数。

我希望能够在 GPU 上执行此操作，并且我已尝试使用 cuBlas。

我的问题是我的矩阵和向量相当小，100x100 矩阵和 100 向量。与 CPU 相比，CuBlas 需要很长时间，我明白为什么，cpu 上的快速缓存和对 GPU 的调用的大量开销。

因此，我试图找出一种聪明的方法来测量将调用传达给 GPU 所需的时间。

这是 CUDA 设置调用并将其发送到图形处理器所需的时间 - 不计算执行矩阵向量乘法实际所需的时间。

我该怎么做呢？

【问题讨论】：

这样您就可以动态选择是将作业发送到 CUDA，还是仅仅出于兴趣？
@Rup ：只是想弄清楚通话的实际成本。并确定慢代码是我的错还是只是架构的产物;-)
对于少量数据，不仅开销会伤害您，而且可能缺乏并行性。 GPU 很大程度上依赖于有足够的线程来隐藏延迟（在 gpu 上比在 cpu 上要糟糕得多）。即使没有调用开销，GPU 也可能比 cpu 慢，除非工作被划分为 lots 个线程。很多很容易就意味着数千个线程。
@Grizzly 我知道 GPU 需要大量线程才能隐藏内存访问。但这引出了一个问题，有多少是很多？ 100 会很多还是我们必须在 1000 或数百万的范围内？
@MartinKristiansen：这取决于需要隐藏什么样的延迟。最佳入住率通常在数万范围内。我通常会说任何少于 1000 个线程的东西可能不会从 gpu 计算中受益，少于 10000 个线程仍然会浪费大部分 gpu 潜力。当然这些都是经验法则，实际上它取决于内核（特别是全局内存访问量，很难隐藏）和使用的 gpu

标签： c++ cuda opencl gpgpu timing

【解决方案1】：

更新：以下结果是针对 2005 硬件（nVidia 7800 GTX）上的手写 FFT GPU 算法，但显示了 CPU-GPU 传输瓶颈的原理

开销不是调用本身，而是 GPU 程序的编译以及 GPU 和主机之间的数据传输。 CPU 针对可完全在缓存中执行的功能进行了高度优化，DDR3 内存的延迟远低于为 GPU 提供服务的 PCI-Express 总线。我在编写 GPU FFT 例程（在 CUDA 之前）时亲身经历过这种情况。请参阅this related question。

N FFTw (ms) GPUFFT (ms) GPUFFT MFLOPS GPUFFT 加速 8 0 0.06 3.352705 0.006881 16 0.001 0.065 7.882117 0.010217 32 0.001 0.075 17.10887 0.014695 64 0.002 0.085 36.080118 0.026744 128 0.004 0.093 76.724324 0.040122 256 0.007 0.107 153.739856 0.066754 512 0.015 0.115 320.200892 0.134614 1024 0.034 0.125 657.735381 0.270512 2048 0.076 0.156 1155.151507 0.484331 4096 0.173 0.215 1834.212989 0.804558 8192 0.483 0.32 2664.042421 1.510011 16384 1.363 0.605 3035.4551 2.255411 32768 3.168 1.14 3450.455808 2.780041 65536 8.694 2.464 3404.628083 3.528726 131072 15.363 5.027 3545.850483 3.05604 262144 33.223 12.513 3016.885246 2.655183 524288 72.918 25.879 3079.443664 2.817667 1048576 173.043 76.537 2192.056517 2.260904 2097152 331.553 157.427 2238.01491 2.106081 4194304 801.544 430.518 1715.573229 1.861814

上表显示了基于内核大小的 GPU FFT 实现与 CPU 实现的时序。对于较小的尺寸，与 GPU 之间的数据传输占主导地位。较小的内核可以在 CPU 上执行，一些实现/大小完全在缓存中。这使得 CPU 成为小型操作的最佳选择。

另一方面，如果您需要对数据执行大批量的工作，而与 GPU 之间的移动最少，那么 GPU 将击败 CPU。

就测量您示例中的效果而言，我建议您进行类似上述的实验。尝试计算为每个矩阵大小计算的 FLOPS，并在 CPU 和 GPU 上针对不同大小的矩阵运行测试。将 GPU 与 CPU 的大小、时间和 FLOPS 输出到 CSV 文件。对于任何分析，请确保您运行代码的数百次迭代并对整个事情进行计时，然后将总时间除以迭代以获得循环时间。如果您的算法允许（例如 10x100 而不是 100x10），也可以尝试不同形状的矩阵。

使用这些数据，您可以了解间接费用是多少。要找出完全一样的实验，但将在 GPU 上执行的内部着色器代码替换为无操作（只需从输入复制到输出）。

希望这会有所帮助，

【讨论】：

【解决方案2】：

您可以通过对缓冲区传输事件使用 clGetEventProfilingInfo 从设备获取事件排队、提交、启动和完成的时间（以纳秒为单位）。

更多信息，以及如何在此处设置：http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/clGetEventProfilingInfo.html

我认为对于 100x100 矩阵，您最好坚持使用 cpu 进行运算。除非您同时有许多要乘法，否则由于（小）传输开销和通常低得多的时钟速度，gpu 的好处几乎不会被注意到。确保调整内核以尽可能多地使用本地数据——在我的硬件上，每个工作组有 32KB，这应该足以容纳两个 100x100 矩阵。内置的点积函数也应该很方便。

去年在 ADFS 上有一个很棒的讨论（见 sessionId: 2908） http://developer.amd.com/afds/pages/OLD/sessions.aspx 他们详细讨论了优化内核和硬编码最佳大小。

【讨论】：

【解决方案3】：

您的矩阵是否已经在 GPU 上？如果没有，CUBLAS 可能会为您传输它们（称为 thunking），这是额外的开销。

此外，GPU 并不能真正用于如此小的计算，即它可能会比 CPU 慢，因为您必须将结果传回。如果可以，请使用更大的矩阵。否则，您可能希望使用流 (cudaStream_t) 在 GPU 上启动多个并行计算。

如果您想在 CUDA 中测量内核的执行时间，则需要将其（或在 GPU 上计算的任何其他内容）包含在事件中，例如使用 CUDA 运行时 API 时：

cudaEvent_t start, stop;

cudaEventRecord(&start);

struct timeval cpuStart, cpuEnd;

gettimeofday(&cpuStart, 0); // get start time on CPU

// Do something with CUDA on the GPU, e.g. call kernels, transfer memory, ...

gettimeofday(&cpuEnd, 0); // get end time on CPU

double seconds = cpuEnd.tv_sec - cpuStart.tv_sec;
double microseconds = cpuEnd.tv_usec - cpuStart.tv_usec;
double cpuDuration = (seconds * 1.0e6 + microseconds) / 1.0e3; // in milliseconds

cudaEventRecord(&stop);

// Wait until the stop event occurred
cudaError_t eventResult;

do
{
  eventResult = cudaEventQuery(stop);
}
while (eventResult == cudaErrorNotReady);

// Assert there was no error; check the CUDA Toolkit Reference for further info
assert(cudaSuccess == eventResult); // requires #include <assert.h> or <cassert>

// Retrieve the time
float gpuDuration = 0.0; // in milliseconds
cudaEventElapsedTime(&gpuDuration, start, stop);

// Release the event objects
cudaEventDestroy(stop);
cudaEventDestroy(start);

您可能需要检查每次调用 CUDA 的错误代码（至少使用断言），因为您可能会从以前的调用中得到错误，从而导致数小时的调试...

（注意：我主要使用 CUDA 驱动程序 API，所以这可能无法开箱即用。抱歉。）

编辑：刚刚看到您想测量调用本身，而不是内核的持续时间。您可以通过简单地测量调用的 CPU 时间来做到这一点 - 请参阅上面的更新代码。这仅适用于 Linux，因为 gettimeofday 不适用于 Windows (AFAIK)。

【讨论】：

在 Windows 上，您可以使用 QueryPerformanceCounter 或 GetSystemTime 等。
设备上的数据我都搞定了，只需要简单的Ax->y，然后把y保存在设备上。
在这种情况下，您可以通过在您的 cublasDgemm() 调用周围放置 gettimeofday()（或 Windows 上的类似方法）来测量 CUBLAS 启动实际内核所需的时间。虽然我自己没有尝试过，但您可以考虑使用 Parallel Nsight（在 Windows 上）或 Visual Compute Profiler（包含在 Linux 上的工具包中）。我现在找不到它，但我确定我已经看到了一些关于 CUDA 4 中的分析钩子的东西......编辑：找到这个 PDF，其中包含一些关于分析 CUDA 的有趣信息：bit.ly/zn6jbP

【解决方案4】：

要查找调用开销，请调用尽可能少的 CUDA 内核。

for (int i=0; i<NLoops; i++) {
    gettimeofday(&cpuStart, 0); // get start time on CPU  

    // Call minimal CUDA kernel  

    gettimeofday(&cpuEnd, 0); // get end time on CPU 

    // save elapsed time
}

按照上面 Alex P. 的代码。

你在内核中做的处理越少，时间差就越大，只是调用开销。

做一些实验来为 Nloops 找到一个合适的值（可能是 1,000,000）。确保经过的时间长于计时器的间隔，否则您最终会得到全零。如果发生这种情况，请编写一些内核代码，在您可以预测的固定时间间隔内执行：（n 个循环，每个循环 x 个）。

很难消除 cpuStart 和 cpuEnd 之间可能发生的所有非 CUDA 计算（如中断处理），但进行多次运行和平均可以产生良好的结果。

【讨论】：