CUDA-非空流中的阻塞流


.主机上非空流是异步流,其上所有的操作都不会阻塞主机执行。相应地,隐式的空流是同步流,大多数添加到空流上的操作都会导致主机在先前所有的操作产生阻塞。

.虽然非空流上在主机上是非阻塞的,但非空流内的操作可以被空流中操作所阻塞。因此可将非空流分为:阻塞和非阻塞两种。

  如果非空流是阻塞流,则空流可以阻塞该非空流中的操作。

  如果非空流是非阻塞流,则它不会阻塞空流中的操作。

三,cudaStreamCreate创建的流是阻塞流,这就意味着流中操作可被阻塞,直到空流中的操作完成,直到空流的执行结束。

流是隐式的同步流,在相同的CUDA上下文情况下,它会和其他所有的阻塞流同步。

.(1)当操作被发布到空流中时,在操作被执行之前,CUDA上下文会等待所有先前操作发布到所有的阻塞流中。

  (2)当操作被发布到阻塞流中时,在操作执行之前,都会被挂起,直至空流中的操作完成。


示例:

(1)阻塞流的阻塞行为:

	const int N =3;
	cudaStream_t *stream = (cudaStream_t*)malloc(sizeof(cudaStream_t)*N);
	for (int i = 0; i < N; i++)
	{
	cudaStreamCreate(&stream[i]);//阻塞流的初始化
	}
	kernel1 << <1,1,0,stream[0]>> > (); //非空流+阻塞流
	kernel2 << <1, 1>> > ();	    //空流
	kernel3 << <1, 1, 0, stream[1] >> > ();//非空流+阻塞流
	for (int i = 0; i < N; i++)
	{
		cudaStreamSynchronize(stream[i]);
	}

执行顺序, kernl2只有在Kernel1执行结束后, kernl3只有在kernel2执行结束后,才会执行。非空流中的阻塞流与空流产生阻塞行。当然,从主机的角度来看,所有kernel都是异步且非阻塞的。 可视化性能分析器(nvvp)如图:


(2)阻塞流的非阻塞行为:

为了消除非空流对空流的阻塞行为,可采用cudaStreamCreateWithFlags(&stream[i], unsiged int flag)进行初始化,flag可取值为[cudaStreamDefault,cudaStreamNonBlocking];前者阻塞与cudaCtreate一样,后者可使得非空流的阻塞行为消失。代码和性能分析结果如图:

const int N = 3;
	cudaStream_t *stream = (cudaStream_t*)malloc(sizeof(cudaStream_t)*N);

	for (int i = 0; i < N; i++)
	{
		cudaStreamCreateWithFlags(&stream[i], cudaStreamNonBlocking);
	}

	kernel1 << <1,1,0,stream[0]>> > ();
	kernel2 << <1, 1>> > ();
	kernel3 << <1, 1, 0, stream[1] >> > ();
	for (int i = 0; i < N; i++)
	{
		cudaStreamSynchronize(stream[i]);
	}

 从图中可以看出,非空流与空流之间并不会产生阻塞影响,都无需等待其他核函数的执行结束。
  • 3
    点赞
  • 22
    收藏
    觉得还不错? 一键收藏
  • 5
    评论

“相关推荐”对你有帮助么?

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论 5
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值