一个OpenCL或CUDA调用的开销是什么?

4
我正在编写一个函数,其中包含大量 BLAS gemv 操作。
我希望能够在GPU上完成这个操作,并尝试使用cuBlas。
我的问题是,我的矩阵和向量非常小,只有100x100的矩阵和100的向量。与CPU相比,cuBlas需要很长时间,我知道为什么,因为CPU具有快速缓存以及调用GPU时的大量开销。
因此,我正在尝试找出一种聪明的方法来测量将调用发送到GPU所需的时间。
也就是说,CUDA设置调用并将其发送到图形处理器所需的时间 -- 不计算矩阵-向量乘法所需的时间。
我该如何做到这一点?

这样你就可以动态选择是将作业发送到CUDA,还是仅出于兴趣的缘故? - Rup
@Rup:只是想弄清楚这个调用到底花费了多少成本,并确定慢代码是我的问题还是架构的产物;-) - Martin Kristiansen
1
对于小量数据,不仅开销会影响性能,而且可能无法实现并行处理。GPU 需要足够的线程来隐藏延迟(在 GPU 上比 CPU 更加严重)。即使没有调用开销,除非将工作分成大量线程,否则 GPU 的速度也可能比 CPU 慢。这里的“大量”通常意味着数千个线程。 - Grizzly
@Grizzly 我知道GPU需要很多线程来隐藏内存访问。但这引出了一个问题,多少才算很多?100个算很多吗?还是我们必须达到1000或百万级别? - Martin Kristiansen
1
@MartinKristiansen:这取决于需要隐藏哪种延迟。最佳占用率通常在数万的范围内。我通常会说,任何少于1000个线程的东西可能不会从GPU计算中获益,而少于10000个线程仍将浪费大部分GPU的潜力。当然,这些都是经验法则,实际上取决于内核(特别是全局内存访问量,这很难隐藏)和使用的GPU。 - Grizzly
4个回答

8
< p >< em >< strong >更新:以下结果是基于2005年硬件(nVidia 7800 GTX)手写的FFT GPU算法,但显示了CPU-GPU传输瓶颈的原理< / em >< / p >

开销不是调用本身,而是GPU程序的编译和在GPU和主机之间传输数据。 CPU高度优化了可以完全在高速缓存中执行的功能,DDR3内存的延迟远低于服务于GPU的PCI-Express总线。在编写GPU FFT例程(CUDA之前)时,我亲身经历过这一点。请参见此相关问题

N       FFTw (ms)   GPUFFT (ms)     GPUFFT MFLOPS   GPUFFT Speedup
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上执行的着色器代码替换为空操作(只需从输入复制到输出)。
希望这有所帮助,

1

为了找到调用开销,调用一个尽可能少的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)。确保经过的时间比计时器的间隔时间长,否则您将得到全部为零的结果。如果发生这种情况,请编写一些在您可以预测的固定时间间隔内执行的内核代码:(x个周期的n次循环)。

很难消除在cpuStart和cpuEnd之间可能发生的所有非CUDA计算(例如中断处理),但是多次运行和平均化可以得到良好的结果。


1

通过在缓冲区传输事件上使用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 他们详细讲解了如何优化内核,并硬编码最佳大小。


1

你的矩阵已经在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调用的错误代码(至少使用assert),因为你可能会从之前的调用中得到错误,导致数小时的调试...

(注意:我主要使用CUDA驱动程序API,所以这可能不适用于开箱即用。对此我很抱歉。)

编辑:刚刚看到你想要测量调用本身,而不是内核的持续时间。 你可以通过在CPU上简单地测量调用的时间来实现这一点-请参见上面更新的代码。 这仅适用于Linux,因为Windows没有gettimeofday(据我所知)。


1
在Windows上,您可以使用QueryPerformanceCounterGetSystemTime等函数。 - Rup
我已经在设备上获取了所有数据,只需要进行简单的Ax->y操作,然后将y保留在设备上。 - Martin Kristiansen
1
在这种情况下,您可以通过在cublasDgemm()调用周围放置gettimeofday()(或Windows上的类似方法)来测量CUBLAS启动实际内核所需的时间。 虽然我自己没有尝试过,但您可以尝试使用Parallel Nsight(在Windows上)或Visual Compute Profiler(包含在Linux工具包中)。 我现在找不到它了,但我肯定看到CUDA 4中有关于分析挂钩的一些信息... 编辑:找到了这个PDF,其中包含有关CUDA分析的一些有趣信息:http://bit.ly/zn6jbP - Alex P.

网页内容由stack overflow 提供, 点击上面的
可以查看英文原文,
原文链接