CUDA 5.0 中cudaStreamCreateWithFlags 的用法

在CUDA 5.0 release note中提到:从CUDA5.0开始,开发者已经可以通过使用cudaStreamCreateWithFlags函数创建可与默认流(null stream)同步执行的cuda stream了。 具体来说,默认流中的Kernel可以与在它后面调用的其他cuda stream中的Kernel同步执行。
我们下面用实际例子说明。(附件中有完整代码)
1、在此例中kernel0为默认流中的kernel;Kernel1为stream[0]中的kernel;Kernel2为stream[2]中的kernel,其中stream[0]与stream[1]为用cudaStreamCreateWithFlags创建的流。流的创建与kernel调用如下:

 /* streamCreate */
  for(int i = 0; i< 2; ++i)
  {
//      cudaStreamCreate(&stream[i]);
   cudaStreamCreateWithFlags(&stream[i], cudaStreamNonBlocking);
  }
  /* kernel launch */
用nvvp进行测试,time line如下图所示:

[attach]2885[/attach] 可以看到处于默认流的kernel0与其他两个kernel同步执行了。

2、下面我们将cudaStreamCreateWithFlags注释掉,换成CUDA 5.0以前使用的cudaStreamCteate。

 for(int i = 0; i< 2; ++i)
  {
cudaStreamCreate(&stream[i]);
//cudaStreamCreateWithFlags(&stream[i], cudaStreamNonBlocking);
  }
用nvvp进行测试,time line如下图所示:

[attach]2883[/attach]
可以看到kernel1与kernel2需要等待kernel0执行完成才可以进行执行。

3、cudaStreamCreateWithFlags中有两个参数第一个参数用于指向待创建的流,第二个参数用于控制创建出的流是否可以与默认流同步执行。当我们将cudaStreamCreateWithFlags中传入的标志cudaStreamNonBlocking修改为cudaStreamDefault后:

  for(int i = 0; i< 2; ++i)
  {
//      cudaStreamCreate(&stream[i]);
   cudaStreamCreateWithFlags(&stream[i], cudaStreamDefault);
  }
nvvp运行结果与前面使用cudaStreamCteate创建流的结果相同,默认流中的Kernel不能与其他流中的Kernel同时执行。

[attach]2884[/attach]

4、值得注意的是,即使使用了cudaStreamCreateWithFlags与cudaStreamNonBlocking标志创建流。如果将被创建出的流中的Kernel放在默认流前面调用,同样不能与默认流同时执行。

 kernel1<<<size_of_array,1,0,stream[0]>>>(d_a,d_c,100);
  kernel2<<<size_of_array,1,0,stream[1]>>>(d_a,d_c,100);
  kernel0<<<size_of_array,1>>>(d_a,d_c,100);
 nvvp运行结果显示,kernel0等到kernel1与kernel2执行完成之后才开始执行。

[attach]2886[/attach]

抢占zehuanwang的原厂沙发,感谢带来一篇详尽的教程,赞!

呵呵,谢谢斑竹支持!

关于为什么0流中的kernel调用必须在其它流之前,我猜测是因为要和以前的cuda版本行为一致.在以前版本的cuda中,0流是最后执行的,这可以用来保证全局存储器的一致性.如果0号流中的kernel在最后调用依旧能够和其它的并行的话,那么它的代码在cuda5上可能产生错误的结果,因此cuda5没有支持这点.

谢谢,虽然还有些不理解的地方

增长知识了,谢谢楼主

增长知识了,谢谢楼主

请问需要包含哪些头文件和静态库,怎么我自己做了一个vs2010的工程,出现了Unresolved extern function 'cudaStreamCreateWithFlags’的错误?

您可能没有使用5.0的库文件,

此函数不需要特别的头文件和库,但您应该至少链接5.0的cudart.lib。

感谢来访。

先谢版主,完整的说我的问题是这样的,我把cuda 5.0 sdk里面的动态并行版的快排(Advanced Quicksort )拿出来研究,发现将vs2010中的工程属性Generate Relocatable Device Code 改成No之后就出现了上面这个错误。然后我自己之前有个工程,原本这里是No,但是我将它改成“-rdc=true”之后,就出现了
1>nvlink : error : Multiple definitions of ‘cudaGetLastError’
1>nvlink : error : Multiple definitions of ‘cudaGetParameterBuffer’
1>nvlink : error : Multiple definitions of ‘cudaLaunchDevice’
1>nvlink : error : Multiple definitions of ‘cudaDeviceSynchronize’
1>nvlink : error : Multiple definitions of ‘cudaDeviceGetLimit’

这样的错误,所以我就有点不明白了,版主所说的这个separation compilation到底应该怎么用?是不是头文件的包含有什么不对,或者有什么地方设置不对?希望得到版主的指导,CUDA新手望见谅~~~

楼主您好,

按理说你是否使用rdc和是否能用cudaStreamCreateWithFlags()无关的。
cudaStreamCreateWithFlags()是个普通的cuda runtime api函数,应该总是能用的(当您连接了cudart.lib后)。

我不明白该项目(sdk examples里的例子), 为何在您禁用rdc后会导致您无法使用cudaStreamCreateWithFlags()。

也许原厂支持可以告诉您,为何该项目禁用rdc会导致提示cudaStreamCreateWithFlags在连接的时候找不到。

感谢周末的深夜来访。

突然想到一点,

如果您不是在host code, 而是在kernel内部使用的cudaStreamCreateWithFlags(), 请您注意连接cudadevrt.lib, 否则将会提示您找不到cudaStreamCreateWithFlags的。

如果您是常规的在host code里调用的cudaStreamCreateWithFlags, 请无视此建议。

感谢周末的深夜来访。

版主是对的,后面重复定义的问题我自己发现了,原来是我在linker和CUDA linker里面都包含了cudadevrt.lib的缘故,谢谢版主的及时解答,祝版主周末愉快