最近尝试用流进行加速,为什么我运行后没有效果?代码如下:
//copy to device
运行结果:
[attach]2903[/attach]
用流后就好像复制执行复制串行了,12.15+6.28+11.96=30.39。
为什么会这样,谢谢大家!
最近尝试用流进行加速,为什么我运行后没有效果?代码如下:
//copy to device
运行结果:
[attach]2903[/attach]
用流后就好像复制执行复制串行了,12.15+6.28+11.96=30.39。
为什么会这样,谢谢大家!
代码如下:
//copy to device
cudaMemcpyAsync(d_a + 0 * n / nstreams, a + 0 * n / nstreams, nbytes / nstreams, cudaMemcpyHostToDevice, streams[0]);
cudaMemcpyAsync(d_a + 1 * n / nstreams, a + 1 * n / nstreams, nbytes / nstreams, cudaMemcpyHostToDevice, streams[1]);
cudaMemcpyAsync(d_a + 2 * n / nstreams, a + 2 * n / nstreams, nbytes / nstreams, cudaMemcpyHostToDevice, streams[2]);
cudaMemcpyAsync(d_a + 3 * n / nstreams, a + 3 * n / nstreams, nbytes / nstreams, cudaMemcpyHostToDevice, streams[3]);
//kernel
init_array<<<blocks, threads, 0, streams[0]>>>(d_a + 0 * n / nstreams, d_c, niterations);
init_array<<<blocks, threads, 0, streams[1]>>>(d_a + 1 * n / nstreams, d_c, niterations);
init_array<<<blocks, threads, 0, streams[2]>>>(d_a + 2 * n / nstreams, d_c, niterations);
init_array<<<blocks, threads, 0, streams[3]>>>(d_a + 3 * n / nstreams, d_c, niterations);
//copy to host
cudaMemcpyAsync(a + 0 * n / nstreams, d_a + 0 * n / nstreams, nbytes / nstreams, cudaMemcpyDeviceToHost, streams[0]);
cudaMemcpyAsync(a + 1 * n / nstreams, d_a + 1 * n / nstreams, nbytes / nstreams, cudaMemcpyDeviceToHost, streams[1]);
cudaMemcpyAsync(a + 2 * n / nstreams, d_a + 2 * n / nstreams, nbytes / nstreams, cudaMemcpyDeviceToHost, streams[2]);
cudaMemcpyAsync(a + 3 * n / nstreams, d_a + 3 * n / nstreams, nbytes / nstreams, cudaMemcpyDeviceToHost, streams[3]);
是的,你这种代码的写法,内存复制是按顺序一个一个拷贝完的,而且,能够overlap的部分在于stream[3]从主机拷贝到设备时,stream[0],stream[1],stream[2]的内核在这段时间内能按顺序执行(假设stream[3]内存拷贝时间够长)。
不妨换个写法试试:
for(int i=0; i<4; i++)
{
cudaMemcpyAsync(xx,xx,xx, HTD, stream[i]);
kernel<<<g,b,0,stream[i]>>(...);
cudaMemcpyAsync(xx,xx,xx, DTH, stream[i]);
}
请在编译的时候注意指定计算能力为2.0或以上,这样的话可以开启HTD和DTH的内存拷贝并发。至于为什么,还请楼主仔细查阅programming gude第三章里面有关concurrent data transfers和overlap behavior的内容!
你好,我那种写法是借鉴《GPU高性能编程CUDA实战》上的,如果用你这种方法的话,理论上会产生阻塞,即后一个HTD必须在前一个DTH后执行,用这种方法测试结果如下,速度慢了2ms。[attach]2905[/attach]
是不是使用流需要对电脑进行额外的配置?还是其他的原因?
我先看看programming guide吧。
我不得不说点,首先是,田园的方法没有错误,也是推荐的写法,楼主不要乱想。
其次,你的结果表明了田园的写法更慢,不一定是田园的错误,因为你这里面还有其他的内容/代码不属于田园。你不能结果不好,就责备我们的版主,请考虑下自己的代码等原因。
最后说一下,你的代码似乎改写自一个例子,该例子里面试图演示用多个流将kernel计算和PCI-E传输同时进行。
(1)对于这个目的,田园的写法是可以的!没有错误!
(2)对于这个目的,如果想表现出节省了总时间的效果,需要有足有的硬件支持(例如双向传输支持),以及需要你的kernel需要稍微长一点,不能传输都20多ms了,计算才6ms, 怎么有效的计算/传输相互掩盖啊!
(3)不建议你用其他的第三方书籍中的做法,指责我们的田园版主的正确做法。
[
请注意我说的,请在编译的时候使用计算能力2.0及其以上!
也就是需要nvcc *.cu -gencode arch=compute_20,code=sm_20
这样才能开启concurret data transfer,使得你的程序能够实现HTD和DTH同时执行!
他这个不妨碍的,因为这个和是否为2.x编译无关。主要是他测试的行为无法展示出执行和复制互相掩盖,以及他的不对的理论理解“如果用田园这种方法的话,理论上会产生阻塞”。
[
这个…concurrent data tranfers好像是需要2.x的设备才支持,似乎还提到了要用page locked的方式分配主存!具体我也没试过~只是记得programming guide里提到了!他们看的书只是讲了一种overlap behavior
我主要推荐他隐藏内存拷贝时间,呵呵!
这么说没错。2个方向上的PCI-E传输需要有2个copy engine的2.0+显卡。
1.x是不行的,GTX460也是不行的。
C2050之类的,和部分GT430,有2个copy engine, 则是可以的。
此外,楼主的卡不行虽然(GT630), 但是楼主弄的好是可以观察到SM上的kernel执行和copy engine的PCI-E数据传输在同时进行,彼此掩盖,节省时间的。
至于上文是否需要为2.x编译,这个无需,即使编译的是1.x的cubin和老的ptx, 依然可以被显卡驱动将ptx即时翻译为当前CPU可执行的二进制码的。
或者简单的说,进程里存在老版本的cubin,不会被执行,也不会破坏掉显卡硬件的dma controller的独立传输能力的。
GTX630支不支持concurrent data transfers我还真不清楚。。。
不过从他执行的结果来看,我哪天打算在M2090上试一下这两种overlap behavior,
从他写的流情况来看,应该是有几个kenel的执行和stream[3]的HTD并发的,但是时间上似乎又不太对!流这个东西我曾经尝试过,也出现过比不用流还慢的情况,真心感觉这东西用起来真是麻烦啊!
嗯嗯。他的结果不怪你,不要听他的。
如同我说过的:“楼主的结果表明了田园的写法更慢,不一定是田园的错误,因为楼主这里面还有其他的内容/代码不属于田园。楼主不能结果不好,就责备我们的版主,请考虑下楼主自己的代码等原因。”
以及他可能的原因,依然在楼上说过,“(2)对于这个目的,如果想表现出节省了总时间的效果,需要有足有的硬件支持(例如双向传输支持),以及需要你的kernel需要稍微长一点,不能传输都20多ms了,计算才6ms, 怎么有效的计算/传输相互掩盖啊!”。
等因素。
streams还是好用,耐用,易用的利器的。
特别是对其他从opencl走来的朋友,不妨考虑下,找到你熟悉的感觉。
刚才查询了programming guide,Concurrent Data Transfer确实是需要2.x或者更高的卡,以及还可能对卡的型号有要求(比如telsa神马的),但是这一功能是H2D(H要求page-locked)copy data的同时,可以执行D2H(H依然要求page-locked)copy。
LZ内容似乎对应于 overlap of Data Transfer and Kernel Execution,这个在1.1或者更高的计算能力的卡上即可实现。
LZ的例子可能改编自SDK里面的例子,或者间接参考自该例子。
强烈推荐LZ看一下CUDA官方博客上的一篇文章,地址是: https冒号//developer.nvidia.com/content/how-overlap-data-transfers-cuda-cc,
这儿不能贴地址,把冒号改成:就行了,你懂的
田园BZ给出的代码对2.X的硬件是最合适的,你给出的代码则不是,这在那篇博文里讲得很清楚。对3.0以上的硬件,这两种代码都可以
BTW,你看的那本书上所用的CUDA硬件是多少的呢?我猜应该是1.X的吧
必须是1.x的,成书之时fermi还没发布呢。
:lol,这就对了。
感谢各位大牛的耐心解答!版主你好!我就一菜鸟,错误之处还请见谅!不过我确实没有指责田园版主的意思。。。
这代码确实来自SDK中的simplestreams,我昨天根据田园版主的建议看了programming guide的对应章节,收益匪浅!,我发现对于我的GPU((deviceProp).asyncEngineCount=1,(deviceProp).concurrentKernels=1)来说,理论上可以支持内存复制核函数执行并行,核函数执行并行,不支持内存复制并行。对于(deviceProp).asyncEngineCount>=2的硬件,田园版主的方法确实是正确的,表示歉意!
另外,
1.这本书里提到了环境变量CUDA_LAUNCH_BLOCKIING的设置,不知道有没有影响。
2.后来我将核函数执行时间增加了,但结果依然不理想,版主能不能提供一个你运行通过的源代码,我想知道是不是我的硬件有问题。
不知道这个compute_20从哪查?对于我的sm_10虽然不支持复制并行,但理论上也支持复制执行并行,执行并行啊,怎么我一点好处也没捞到。
完全正确!