如何客观评估GPU速度（没看到GPU比CPU快）

[toshibahpdopod]

我用hotball给出的GPU测试代码2（second_cuda.cu）在一台HP Z800的工作站（CENTOS5.4 32G内存，TESLA C1060,FX1800显示卡各一块）上编译通过，运行时候发现计算耗时不稳定。
$ ./a.out
Max error: 1.19209e-07 Average error: 4.22521e-08
Time used: 0.1100 (18.18 GFLOPS)
$ ./a.out
Max error: 1.19209e-07 Average error: 4.22521e-08
Time used: 0.0900 (22.22 GFLOPS)
$ ./a.out
Max error: 1.19209e-07 Average error: 4.22521e-08
Time used: 0.1100 (18.18 GFLOPS)
$ ./a.out
Max error: 1.19209e-07 Average error: 4.22521e-08
Time used: 0.0700 (28.57 GFLOPS)
$ ./a.out
Max error: 1.19209e-07 Average error: 4.22521e-08
Time used: 0.0700 (28.57 GFLOPS)

另外编写了简单的cuda程序，发现计时比较，并不比CPU算得快啊
哪位支持一下，推荐个小示范程序，能体现gpu比cpu算得快好吗
谢谢

[cuda2010]

根据我的经验，GPU上程序运行时间通常非常稳定，不可能出现这么大的涨落。性能方面20Gflops左右的成绩也严重偏低，离tesla的理论性能差得太远。最好把完整测试代码贴出来，才能知道是什么问题。

[meppp]

测试性能不稳定可能是由于clock函数不精确造成的，你应该使用更加精确的计时函数，比如performance counter

[zenny_chen]

呵呵，主要问题是出在存储器访问开销。
因为这里面既可能有Cache命中问题，也可能有页缺失异常问题，以及TLB的一些问题。clock的精度问题几乎可以被忽略掉。

[toshibahpdopod]

非常感谢各位指点，我也在查资料，有种说法除了优化程序外，为了测试时间更加精确，应该在计时前先激活一下cuda，比如执行一下cudaFree(NULL) ，第一次激活cuda会耗时长一些

[无心人_过过小日子]

[Quote=引用 11 楼 zenny_chen 的回复:]
呵呵，主要问题是出在存储器访问开销。
因为这里面既可能有Cache命中问题，也可能有页缺失异常问题，以及TLB的一些问题。clock的精度问题几乎可以被忽略掉。
[/Quote]

GPU上有这些东西吗?呵呵.

[cuda2010]

没错，第一遍执行涉及一些cuda初始化的东西，涉及到所谓的warm up问题。
我测试程序的时候一般这样计时：

test<<<xxx,xxx>>>();
cudaThreadSynchronize();
t0=timer();
test<<<xxx,xxx>>>();
cudaThreadSynchronize();
t1=timer();
dt=t1-t0;

[Quote=引用 12 楼 toshibahpdopod 的回复:]
非常感谢各位指点，我也在查资料，有种说法除了优化程序外，为了测试时间更加精确，应该在计时前先激活一下cuda，比如执行一下cudaFree(NULL) ，第一次激活cuda会耗时长一些
[/Quote]

[OpenGPU2010]

学习一下！！

---------------------------------------------------------------------

开源图形处理器体系结构论坛(OpenGPU论坛)
http://www.opengpu.org/bbs/

OpenGPU Graphics Open Source community（ 图形开源社区），聚焦领域(focus domain)包括：
* GPU Architecture（图形处理器体系结构）.
* Graphics Algorithm（图形算法).
* Open Source Rendering Engine（开源渲染器）.
* Open Source GPU Simulator/RTL Implement（开源GPU模拟器）.
* GPGPU Programming （ 面向通用的图形处理器编程）
* GPU General-purposed Computing （ GPU通用计算）.
.

[zenny_chen]

[Quote=引用 13 楼 l7331014 的回复:]

引用 11 楼 zenny_chen 的回复:
呵呵，主要问题是出在存储器访问开销。
因为这里面既可能有Cache命中问题，也可能有页缺失异常问题，以及TLB的一些问题。clock的精度问题几乎可以被忽略掉。


GPU上有这些东西吗?呵呵.
[/Quote]
呵呵，楼主在测时间的时候一定在CPU这边测的吧。

[zenny_chen]

当然，GPU上的一些texture cache之类的也可以考虑进去，呵呵。

[zenny_chen]

现在不管是GPU也好，CPU也罢。
速度已经快的可以了。
现在主要瓶颈就是Memory transferring。而且随着CPU、GPU的发展会显得更加明显。
所以，如何充分利用存储器资源，使得时间不会被白白地流失掉就需要好好设计。
CPU、GPU目前往往都是通过提高吞吐来回避这个问题。对于CPU而言，偶使用SSE对128位数据进行读取，这个时候数据可能会从外部存储器被放入L2或L3，利用这段时间做其它计算，使得总线操作和CPU操作能够并行处理。
而GPU这边就是利用其大规模并行计算的优势来进行化解。当某个线程（或者说，对于nVidia的ptx2.0架构而言，一个warp）因为存储器资源问题被阻塞住，那么其它CTA可以继续执行。另外，对于分支所产生的延迟也是通过类似技术来规避的，呵呵。

所以现在在处理器这端出现了L3，确实是挺有意思，而且也带有些无奈。这实际上也是一种分流总线阻塞的方法。
另外AMD有一个叫Transport的技术很有意思。它把总线搞成双工的模式，使得读、写能够同时进行，呵呵。不过具体的效果偶也没做过深入研究。

[OpenGPU2010]

[Quote=引用 18 楼 zenny_chen 的回复:]

现在不管是GPU也好，CPU也罢。
速度已经快的可以了。
现在主要瓶颈就是Memory transferring。而且随着CPU、GPU的发展会显得更加明显。
所以，如何充分利用存储器资源，使得时间不会被白白地流失掉就需要好好设计。
CPU、GPU目前往往都是通过提高吞吐来回避这个问题。对于CPU而言，偶使用SSE对128位数据进行读取，这个时候数据可能会从外部存储器被放入L2或L3，利……
[/Quote]


：〉学习了

---------------------------------------------------------------------

开源图形处理器体系结构论坛(OpenGPU论坛)
http://www.opengpu.org/bbs/

OpenGPU Graphics Open Source community（ 图形开源社区），聚焦领域(focus domain)包括：
* GPU Architecture（图形处理器体系结构）.
* Graphics Algorithm（图形算法).
* Open Source Rendering Engine（开源渲染器）.
* Open Source GPU Simulator/RTL Implement（开源GPU模拟器）.
* GPGPU Programming （ 面向通用的图形处理器编程）
* GPU General-purposed Computing （ GPU通用计算）.
.

[zenny_chen]

呵呵，感谢楼上OpenGPU让偶有机会了解OpenGPU网站。
那个站确实不错，也是一同研究Graphics Architecture的好地方，呵呵。

[zenny_chen]

呵呵，对了。
CUDA2010在OpenGPU里也是很活跃啊，呵呵。

[cuda2010]

对了，关于计时精度，最近观测到在某些极端情况下gpu初始化[显式或隐式调用cudaSetDevice()]的用时可以长达12秒。这个如果放入计时那结果就完全不准了。不知道其他人有没有观察到这种现象。

[lxw886]

[Quote=引用 18 楼 zenny_chen 的回复:]
现在不管是GPU也好，CPU也罢。
速度已经快的可以了。
现在主要瓶颈就是Memory transferring。而且随着CPU、GPU的发展会显得更加明显。
所以，如何充分利用存储器资源，使得时间不会被白白地流失掉就需要好好设计。
CPU、GPU目前往往都是通过提高吞吐来回避这个问题。对于CPU而言，偶使用SSE对128位数据进行读取，这个时候数据可能会从外部存储器被放入L2或L3，利用……
[/Quote]
学习了……，chen是不是研究体系结构的？

[victorypiter]

Mark一下 学习了。。

[thouden]

[Quote=引用 8 楼 toshibahpdopod 的回复:]
非常感谢，另外请教一下，有没有cuda的简单例子程序，能显示出GPU做加减乘法比CPU算得快很多的
[/Quote]
在CUDA的C/src下面有很多的示例程序
您可以尝试运行一下
我跑8192*8192矩阵的二维8192*8192FFT运算 在Windows下需要大概500ms 在linux下运行大概只需要100ms左右 楼主可以尝试下一下

[thouden]

[Quote=引用 22 楼 cuda2010 的回复:]
对了，关于计时精度，最近观测到在某些极端情况下gpu初始化[显式或隐式调用cudaSetDevice()]的用时可以长达12秒。这个如果放入计时那结果就完全不准了。不知道其他人有没有观察到这种现象。
[/Quote]
个人感觉这个时间应该放在初始化里面吧

[taseaa]

[Quote=引用 29 楼 thouden 的回复:]

引用 8 楼 toshibahpdopod 的回复:
非常感谢，另外请教一下，有没有cuda的简单例子程序，能显示出GPU做加减乘法比CPU算得快很多的

在CUDA的C/src下面有很多的示例程序
您可以尝试运行一下
我跑8192*8192矩阵的二维8192*8192FFT运算 在Windows下需要大概500ms 在linux下运行大概只需要100ms左右 楼主可以尝试下一下
[/Quote]

不会吧，五倍的差距哟。

[Cyrosly]

差距那么大不太可能，应该在相同的配置下对比，包括主板，CPU等。

不过GPU和CPU的加速比实际没那么夸张，除非跑那些非常一目了然的并行算法，但实际应用中算法的并行复杂度往往不是那么简单。