上节中描述到CUDA中可以通过流来实现网格级并发，按照流的划分主要分为两种：

空流（默认创建）
非空流（需要使用API创建）

空流为一个同步流，在空流中的大部分操作都会堵塞主机，kernel运行除外，

而非空流为一个异步流，其上所有的操作都不堵塞主机执行。

而非空流按照与空流直接的关系还可分为两种类型

堵塞流
非堵塞流

这个时候可能就有点晕乎，可以借用另外一个博客大牛的图片来说明：

非空流中的堵塞和非堵塞主要是针对与空流之前的关系而划分，下面来根据几种情况来说明：

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>> > ();     //空流 会堵塞kernel 2和kernel 3
kernel2 << <1,1,0,stream[0]>> > (); //非空流+阻塞流
kernel3 << <1, 1, 0, stream[1] >> > ();//非空流+阻塞流
for (int i = 0; i < N; i++)
{cudaStreamSynchronize(stream[i]);
}

上述例子中kernel 1为空流中的操作， kernel2为stream[0], kernel3为stream[1], 上述三个操作对主机而言都是异步的，所有需要进行同步操作，当kernel分别加载到流中，在GPU真正的执行顺序并不是我们想象的三个kernel同时并行，而是先执行kernel1, kernel1执行完后才会同时执行kernel 2和3，这是因为空流对于非空流的操作具有一个堵塞作用。非空流的操作可以被空流中的操作所堵塞，即kerne1执行完后，才会执行kernel 2和3.

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

上句话的意思是当kernel2执行之前，会一直等待kernel 1执行完毕之后才会执行。

2：非空堵塞流对空流行为

非空堵塞流中的操作会堵塞空流中的操作。

将上述用例进行稍微改造：

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>> > ();      //空流 会堵塞kernel 3
kernel3 << <1, 1, 0, stream[1] >> > ();//非空流+阻塞流
for (int i = 0; i < N; i++)
{cudaStreamSynchronize(stream[i]);
}

其中kernel1操作加入到非空流stream[0]中， kernel为一个空流操作，kernel3为一个非空流stream[1]中，由于非空堵塞流中的操作会堵塞空流中的操作，所以kernel2会一直等待kernel1执行完毕之后，才会得到执行，而kernel2为空流，同时会堵塞kernel3,所以执行顺序为kernel1->kernel 2->kernel 3串行执行的效果。

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

这句话的意思是kernel 2在执行之前，会一直等待之前的非空堵塞流操作执行完毕，即kernel1执行完毕才会执行kernel2.

3：非空非堵塞流

非空非堵塞流不会对空流操作堵塞操作，空流也不会堵塞非空非堵塞流。

CUDA中使用cudaStreamCreate ()API创建的是非空堵塞流，如果要创建非空非堵塞流需要使用到如下API:

__host__ __device__ cudaError_t cudaStreamCreateWithFlags ( cudaStream_t* pStream, unsigned int flags )

其中flags为创建的流类型：

cudaStreamDefault：默认值，创建流为非空堵塞流，会对空流造成堵塞。

cudaStreamNonBlocking：创建非空非堵塞流，不会对空流造成堵塞。

下面非一个非空非堵塞流的例子：


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]);
}

上述例子kernel 1和kernel 3为一个非空流， kernel 1和kernel 2,kernel 3为并行执行。

4：非空流之间

非空流之间不管是堵塞还是非堵塞，都不会相互影响，在上章节中的例子：

or (int i= 0; i< nStreams;i++)
{int offset = i * bytesPerStream;cudaMemcpyAsync(&d_a[offset], &a[offset], bytesPerStream, streams[i]);kernel<<grid, block,0,streams[i]>>(&d_a[offset]);cudaMemcpyAsync(&a[offset], &d_a[offset], bytesPerStream,stream[i]);
}for(int i =0; i< nStream;i++)
{cudaStreamSynchronize(streams[i]);
}

即使创建的流为非空流，相互之前也不会受影响。

参考资料

https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__STREAM.html#group__CUDART__STREAM_1gb1e32aff9f59119e4d0a9858991c4ad3

https://blog.csdn.net/qq_17239003/article/details/78994992

CUDA 网格级并发-流(2)相关推荐

CUDA 网格级并发-流(1)
CUDA中的程序并发可以分为两种内核级并发网格级并发内核级并发是开发程序中经常使用到的,即通过划分block和thread实现同一个内核在GPU上同时并发,将同一个核分别部署到不同的SP上进行同 ...
远程办公是巨头游戏？十倍扩容，他们如何做到百万级并发流量
疫情发生后,除了Zoom这样深耕视频会议多年的软件,钉钉.企业微信.飞书等一大批互联网巨头也开通了免费服务,凭借着自身庞大的资源四处招揽用户. 据说,远程办公工具是2020年的第一个风口. 疫情发生后 ...
蚂蚁金服面对亿级并发场景的组件体系设计
来自:蚂蚁金服公众号mPaas 作者:吕丹(凝睇),2011 年加入支付宝,先后负责了支付宝 Wap.alipass 卡券.SYNC 数据同步等项目,并参与了多次双十一.双十二.春节红包大促活动,在客 ...
淘宝从几百到千万级并发的十四次架构演进之路！
点击上方好好学java ,选择星标公众号重磅资讯.干货,第一时间送达今日推荐:用好Java中的枚举,真的没有那么简单!个人原创+1博客:点击前往,查看更多链接:https://segmen ...
Github星标90K？京东架构师一篇讲明白百亿级并发系统架构设计
学习高并发系统设计的原因高并发到底是什么,想必各位多多少少对此都有所了解,那我在这就不多说了.真正经历过"双11"以及"618"的小伙伴应该都知道,在大促时如 ...
如何打造千万级Feed流系统
摘要: Feed流是一个目前非常常见的功能,在众多产品中都有展现,通过Feed流可以把动态实时的传播给订阅者,是用户获取信息流的一种有效方式.在大数据时代,如何打造一个千万级规模的Feed流系统仍然是 ...
家乐福618保卫战二-零售O2O场景中的万级并发交易情况下的极限性能调优
本系列简介这个系列可以帮助普通程序员们深刻的意识到平时工作中到底还有什么不足以及如何进一步进化成真正意义上的架构师.CTO以及后面的道路是如何走的: 这个系列可以帮助企业IT管理者深刻意识到,性能安 ...
EMQ 助力阿里云洛神云网络构建新一代“亿级并发、百万级吞吐”NLB 网络型负载均衡系统
万物智联的数字化时代,我们正走在从"数据量变"到"连接质变"的道路上.在日益丰富的物联网应用场景中,实现海量设备与云端之间双向通信连接,分析并从中获得实时洞察成 ...
阿里巴巴十亿级并发系统设计手册已开源（2022最新版）
高并发,几乎是每个程序员都想拥有的经验.原因很简单:随着流量变大,会遇到各种各样的技术问题,比如接口响应超时.CPU load升高.GC频繁.死锁.大数据量存储等等,这些问题能推动我们在技术深度上不断 ...

CUDA 网格级并发-流(2)