在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]