最近尝试用流进行加速，为什么我运行后没有效果？代码如下：
//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虽然不支持复制并行，但理论上也支持复制执行并行，执行并行啊，怎么我一点好处也没捞到。
完全正确！