CUDA中的程序并发可以分为两种

内核级并发
网格级并发

内核级并发是开发程序中经常使用到的，即通过划分block和thread实现同一个内核在GPU上同时并发，将同一个核分别部署到不同的SP上进行同时运行，上述并发方式为我们经常使用到的方式。

然而在内核级并发中，其实还是存在同步忙等待现象，比如内存copy时，将数据从CPU copy到GPU或者在GPU中的数据结果 copy到CPU中这些过程都需要进行忙等待，以及有时经常遇到两个不相关的内核函数可以同时运行以提升以最大利用率，此时就需要用到CUDA中的网格级并发，网格级并发运行多个内核同时运行在同一个硬件设备中。

流和事件

CUDA中的网格级并发是通过建立不同的流来实现。CUDA流是一系列异步的CUDA操作，这些操作按照主机代码确定的顺序在设备上执行。流通过将这些主机代码的操作进行封装，以保持操作的顺序，并且允许这些操作在流中排队，按照先后顺序执行所有的操作。在相同的流中的操作，按照严格的顺序来执行，而在不同流中的操作相互不影响，所以可以通过建立不同的流，来实现网格级并发。

CUDA中的流分为两种隐式和显式流。

隐式流又称为空流，CUDA在内核其他时会默认创建一个流为隐式流，如果不创建新的流，则所有的操作将封装到默认流中。默认流的对数据传输采用同步，对内核运行采用异步方式。数据传输API一般为cudaMemcpy等，而kernel运行则使用异步方式，这就意味着CUDA 主机端在提交内核之后会立即返回，主机端可以在等待kernel运行期间做其他操作，即可以将主机和设备端的计算这一段过程重叠并行运行。

显式流可以调用CUDA API进行显式的创建，为一个异步流，即封装的所有操作都可以是异步的，包括内存copy，例如

cudaMemcpyAsync(void *dst, const void *srcm, size_t count, cudaMemcpyKind kind, cudaStream_t stream=0)

上述内存copy异步动作必须显式的设置一个流，调用完后即使数据没有copy完毕，也会立即进行返回。

值得注意的一点是，使用cudaMemcpyAsync()函数时，使用的内存必须是固定主机内存，可以使用以下函数进行申请：

cudaMallocHost(void **ptr, size_t size);

cudaHostAlloc(void **pHost, size_t size, unsigned int flags);

在主机虚拟内存中固定分配，可以确保其在CPU内存中的物理位置在应用程序的整个生命周期中保持不变。否则，操作系统可以随时自由改变主机虚拟内存的物理位置。如果在没有固定主机内存的情况下执行一个异步CUDA操作，操作系统可能会在物理层面上移动数组，而CUDA操作运行时将该数组移动到设备中，这样会导致为定义行为。

CUDA相关API

CUDA中有关流的API如下：

创建一个非空流：

__host__ cudaError_t cudaStreamCreate ( cudaStream_t* pStream )

此时创建的流为一个异步流。

销毁一个非空流：

__host__ __device__ cudaError_t cudaStreamDestroy ( cudaStream_t stream )

CUDA程序中一般典型的模式为：

1：将输入数据从主机移动到设备上。

2：在设备上执行一个内核。

3：将结果从设备移回主机中。

在默认流中，对数据在主机和设备之间移动是通过cudaMemcpy()函数来实现，而该函数是一个同步堵塞函数。而在一个非空流中为一个异步流，可以将主机和设备之间数据移动也改为一个异步操作，这样就能隐藏该步骤的时间消耗。

异步copy函数API为：

__host__ __device__ cudaError_t cudaMemcpyAsync ( void* dst, const void* src, size_t count, cudaMemcpyKind kind, cudaStream_t stream = 0 )

下面一个例子说明如何实现网格级并发，多个流同时运行：

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

其中cudaStreamSynchronize()函数为强制堵塞主机，一直等到给定流中所有的操作都完成，函数原型为：

__host__ cudaError_t cudaStreamSynchronize ( cudaStream_t stream )

执行的流过程如下：

与传统的串行方式相比，在性能提升上有很大改进。其中数据传输操作虽然分布在不同的流中，但是由于硬件资源PCIe总线限制，无法真正并行执行，但是可以利用kernel异步执行过程中，来进行下一个流的传输，以此来隐藏传输数据的时间消耗，来提高性能。

重叠GPU和CPU执行

CUDA中还提供了另外一种方法用于判断一个异步流中的所有操作是否完毕cudaStreamQuery(),用于查询所有流中的操作是否已经完成，该函数非堵塞某事，函数原型为：

__host__ cudaError_t cudaStreamQuery ( cudaStream_t stream )

当流中的操作都完成时，函数返回cudaSuccess，当流中仍然有操作在执行时返回cudaErrorNotReady。利用cudaStreamQuery()函数可以完成来实现GPU和CPU同时执行。

下面以一个内核实现向量与标量加法例子来说明重叠GPU和CPU执行是如何实现的。

内核kernel的实现加法运算如下：

__global__ void kernel(float * g_data,float value)
{int idx = blockIdx.x * blockDim.x + threadIdx.x;g_data[idx] = g_data[idx] + value;
}

使用异步操作完成数据copy，计算等操作，如下：

cudaStream_t Stream;
cudaStreamCreate (&Stream);
cudaMemcpyAsync(d_a,h_a,nbytes,cudaMemcpyHostToDevice, Stream);
kernel<<<grid, block,Stream>>>(d_a, value);
cudaMemcpyAsync(h_a, d_a,nbytes, cudaMemcpyDeviceToHost, Stream);

异步等待时刻,使用cudaStreamQuery()不断查询是否流中的操作是否已经完毕：

unsigned long int counter = 0;
while（cudaErrorNotReady == cudaStreamQuery()）
{counter++
}

在循环等待过程中，CPU不断执行counter++操作，以实现GPU和CPU同时运算。

参考资料

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

《CUDA C编程权威指南》

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

CUDA 网格级并发-流(2)
上节中描述到CUDA中可以通过流来实现网格级并发,按照流的划分主要分为两种: 空流(默认创建) 非空流(需要使用API创建) 空流为一个同步流,在空流中的大部分操作都会堵塞主机,kernel运行除外, ...
远程办公是巨头游戏？十倍扩容，他们如何做到百万级并发流量
疫情发生后,除了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 网格级并发-流(1)

流和事件

CUDA相关API

重叠GPU和CPU执行

参考资料

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

最新文章

热门文章