我看手册里面说，采用异步+流的话可以使一部分数据复制和计算同时进行，如果硬件支持concurrent kernel还可也使不同流中的kernel同时进行计算。以下是我的问题：
1.如果不采用流，只用异步，如果先进行数据拷贝后进行计算，那么设备还没有完成复制前就被返回给主机线程，然后呢？主机线程会在复制还没进行完的时候调用kernel吗？
2.有段程序实现如下:
for(int i=0; i<n; i++)
{
for(int j=1; j<m[i]; j++)
{
kernel<<<grid_dim, blk_dim, 0 , streams[i]>>>(pout[i][j], pin[i],pout[i][j-1]);
}
}
我解释一下，之所以采用双重循环有2个原因，首先是逻辑上对i为某固定值时，j+1步依赖于j步的计算结果，所以内循环必须要串行，然后，对不同的i，内循环的次数不同，即当i1 != i2 时， m[i1] != m[i2]。但是呢，不同的i之间，是没有数据依赖的。外循环反而是可以并行的。
如果只采用cuda流以实现concurrent kernel的话，感觉比较困难，不直接。所以我想采用最简单的openmp来展开外循环，不知道这样是否可以？
还有就是如果采用流，kernel函数好像就必然是异步运行的，在内循环里面j+1步存在对j步的数据依赖的情况下，能否保证j步完成之后才进行j+1步呢？
LZ您好，您的问题答复如下：
1：“如果不采用流，只用异步，如果先进行数据拷贝后进行计算，那么设备还没有完成复制前就被返回给主机线程，然后呢？主机线程会在复制还没进行完的时候调用kernel吗？”
如果您的异步copy和kernel调用在同一个stream里面，那么是串行的，依然会保持前一步完成之后，后一步才执行，此时是安全的。
如果您的异步copy和kernel调用在不同的两个stream里面（都不在默认stream），那么是不保证顺序的，但是此时不同stream里的操作，原本就应该保证是不相关的，所以也不会有问题。
如果您异步copy之后，直接不加stream参数地启动kernel，此时kernel在默认stream里面，会产生一次隐式的全局同步。（除非您指定异步copy的流和默认流不进行同步）
2：您的外循环每次循环是没有依赖的，那么的确可以用OPENMP展开。
2.1：您的内循环直接循环的话，如果内循发布kernel的时候总是在同一个stream里面，kernel总是串行执行的，并无问题。
以及补充的2.2：稍微说一下“异步”不同的含义。当我们谈及kernel启动是异步操作，此时的含义是，该操作会在执行结束之前返回；而如果谈及不同stream的操作是异步的，指的是不同stream之间并不保证先后顺序。
所以，在您最后一段中，如果您的“异步”是上述第一种含义，那么是无需强调的，以及您需要将循环发射的kernel放到同一个stream里面才能保证j+1步对j步的数据依赖。（以及您的代码表示您实际上是内循环的kernel都在同一个stream里面）
如果您的“异步”是第二种含义，此时您是将kernel发射到不同的stream里面了，此时是无法保证j+1步对j步的数据依赖的。（但是您的代码表示您不是这种情况，因此您的文字叙述中对第一种“异步”含义下的代码实现，担心第二种“异步”可能造成的问题，是没有必要的。）
大致如上，供您参考。
祝您好运~
增加了一些补充说明内容。
继续追问，我说的异步的确是第一种意思：该操作会在执行结束之前返回。就是说，如果是在同一个流里面的话，虽然是用了异步的函数，虽然提前返回了，但是会在下一个函数执行之前进行隐式同步是吧？
LZ您好：
1：一个流里面，即使是异步的函数，也会顺序执行的。比如你连续发射多个kernel，他们都是异步的，但是都会保证一个执行完毕，下一个才上。
2：上述这种行为一般不用“隐式同步”形容。上文中提及的隐式同步指的是默认流（stream 0）的一些操作会导致其他流都停下来等待。
大致如此，祝您好运~
谢谢，我明白了，顺便问一句，kepler在异步与流这方面有啥改进（看他的宣传ppt好像更nb的样子）
LZ您好，计算能力3.5的kepler GPU具备Hyper Q功能，能够更好地识别代码中多个流之间的依赖关系并发掘其并行性。这包括copy和计算并行以及多个kernel同时计算。
在之前的GPU上，可能时常需要手工调整为GPU能识别的等价写法，才能有较好的的效果，但在SM3.5的kepler上，这种对书写的要求降低了，GPU硬件能直接识别，这减少了调试代码的工作量。
具体您可以参考Hyper Q技术的相关说明和示例。
以及SM3.5的kepler还有一个功能是允许kernel自己调用kernel，但这个和您问的关系不大，就不展开了。
大致如此，祝您编码顺利~