请指导如何提高性能

加速计算 CUDA
1

那边发了没人理，发这边，似乎人气旺一些，请大家不吝赐教，谢谢。

我这个程序是基于蒙特卡洛算法的。不过问题不限于此，不了解这个算法的，可以只当有那么一个简单的算法就行了。

首先简单说一下蒙特卡洛算法，给不了解得人，严格地说，是我这里使用的蒙特卡洛算法，仅仅是其算法的一种应用而已。
蒙特卡洛计算是：通过对参数使用正态分布随机数处理，将处理后的结果进行运算。重复以上步骤，并统计运算后的结果出现在某个范围内的概率。
Kernel函数里:


int i=blockIdx.x*blockDim.x+threadIdx.x;//索引

float tro1=sqrtf(-2*__logf(rd1))*__cosf(2*PI*rd2)*105+1050;//rd1，rd2是内容为随机数的数组

float pi1=sqrtf(-2*__logf(rd1))*__cosf(2*PI*rd2)*513.3f*0.1f+513.3f;//以上两行的最后两个数，应该是上游函数传递的参数，为简单起见，以常数代替。

float Sigma=pi1*4.5625f*0.5f;

float L=__fdividef(powf(10,-1000*(27.1399f-4.39786f*__log10f(Sigma))),(tro1+273));

float Ft=1000*L;//每一步都是使用上一步的计算结果，最后得到Ft

if(Ft*1.25f > 1)

  atomicAdd(Ni,1);//Ni为统计变量，统计Ft在给定范围内出现的次数。由于出现的可能性较小，故使用原子函数

Host端主要为建立rd1，rd2两个数组并赋值。启动Kernel，最后返回 Ni/次数（即线程数）。10次运算的平均时间大概是0.36s左右（Tesla C1060），加速比大概也有十几二十倍，记不清了。但我最初设想，因为蒙特卡洛相当于若干线程做同一件事情，最后汇总结果，所以加速比应该接近线程数。（我以前用CPU，循环手动展开，最好时也有4倍左右加速）
试过几个方法，都没有效果：分成多个流，异步传输计算；使用WriteCombined；使用Mapped（这个运行失败，硬件不支持？）
初步考虑是数据量太小？

请各位高人不吝赐教，万分感谢！
（附件，sampleStream的是用流的，sampleWCM的是用WriteCombined，sampleMapped是用Mapped的（这个我这运行不了）），由于测试，都没有用cutilSafeCall，请谅解。

有我说的不清楚的地方，请提出，谢谢

[ 本帖最后由 sailorsb 于 2010-1-13 13:33 编辑 ]

2

首先我想问的是：你的加速比是怎么计算的，是只kernel，还是整体的？你用的是什么硬件？这些都有影响。

不过加速一二十是非常不错的:)

3

加速比是整体的。CPU代码是纯串行的，基本上可以认为就是把GPU所作的工作用循环来做。二者计时都是全部的时间。重复10次计总时间。由于连续运行几次，应该不存在“唤醒”的问题。GPU是Tesla C1060，只用了一个核心。CPU记不清了，y应该是INtel 的E5520。也是只用了一个核心。

我觉得现在的加速比不高，是因为CPU如果循环展开4次或8次的话，也能够有接近4倍的加速比。而且，我觉得这个算法基本上是：做20万次相同的计算（其中个别参数不同），最后统计结果。应该是并行度很高的算法。

原来以为加速比能够接近GPU核心数，但现在差的太多了。:eek:

4

我看过了你的代码，首先声明一下，你的代码是有错误的，至少我下载的版本有一些小错误。另外，你用了分支和原子指令，这个要考虑到。

5

另外数据分配和传输的时候是很长的，我最近也在烦这个。如果你测试kernel的话，可能加速比接近ＧＰＵ核心数！？:slight_smile:

6

我代码中sampleMapped。cu这个文件应该是有错的。其他两个应该都是对的，我实机验证过。三个是单独的，之间没有关系。

没想到数据分配和传输的时间那么长。有机会的话，听你的，测试一下kernel的时间

7

:slight_smile: