CUDA学习之流和并发--part1

最新推荐文章于 2023-03-11 12:18:44 发布

我想静静，

最新推荐文章于 2023-03-11 12:18:44 发布

阅读量1.1k

点赞数 1

分类专栏： CUDA 文章标签： cuda

本文链接：https://blog.csdn.net/weixin_42764932/article/details/115478659

版权

CUDA 专栏收录该内容

29 篇文章 14 订阅

订阅专栏

一般来说，在CUDA C编程中有两个级别的并发：

内核级并发
网格级并发

之前介绍了内核级的并发，在此级别的并发中，单一的任务或内核被GPU的多个线程并行执行。
提升内核性能的几种方法，分别是从编程模型、执行模型和内存模型的角度进行介绍的。

本章将研究网格级的并发。在网格级并发中，多个内核在同一设备上同时执行，这往
往会让设备利用率更好。

6.1 流和事件概述

CUDA流是一系列异步的CUDA操作，这些操作按照主机代码确定的顺序在设备上执行。

这些操作包括：

在主机与设备间进行数据传输
内核启动
以及大多数由主机发起但由设备处理的其他命令。

流能封装这些操作，保持操作的顺序，允许操作在流中排队，并使它们在先前的所有操作之后执行，并且可以查询排队操作的状态。

流中操作的执行相对于主机总是异步的。CUDA运行时决定何时可以在设备上执行操作。我们的任务是使用CUDA的API来确保一个异步操作在运行结果被使用之前可以完成。

在同一个CUDA流中的操作有严格的执行顺序，而在不同CUDA流中的操作在执行顺序上不受限制。

使用多个流同时启动多个内核，可以实现网格级并发。

因为所有在CUDA流中排队的操作都是异步的，所以在主机与设备系统中可以重叠执行其他操作。在同一时间内将流中排队的操作与其他有用的操作一起执行，可以隐藏执行那些操作的开销。

CUDA编程的一个典型模式是以下形式：

将输入数据从主机移到设备上。
在设备上执行一个内核。
将结果从设备移回主机中。

在许多情况下，执行内核比传输数据耗时更多，可以完全隐藏CPU和GPU之间的通信延迟。

通过将内核执行和数据传输调度到不同的流中，这些操作可以重叠，程序的总运行时间将被缩短。

流在CUDA的API调用粒度上可实现流水线或双缓冲技术。

CUDA的API也分为同步和异步的两种：

同步行为的函数会阻塞主机端线程直到其完成
异步行为的函数在调用后会立刻把控制权返还给主机。

异步行为和流是构建网格级并行的支柱。

虽然我们从软件模型上提出了流，网格级并行的概念，但是说来说去我们能用的就那么一个设备，如果设备空闲当然可以同时执行多个核，但是如果设备已经跑满了，那么我们认为并行的指令也必须排队等待——PCIe总线和SM数量是有限的，当他们被完全占用，流是没办法做什么的，除了等待。

我们接下来就要研究多种计算能力的设备上的流是如何运行的。

6.1.1 CUDA流

所有CUDA操作都是在流中进行的，虽然我们可能没发现，但是有我们前面的例子中的指令，内核启动，都是在CUDA流中进行的，只是这种操作是隐式的，所以肯定还有显式的，所以，流分为：

隐式声明的流，我们叫做空流
如果我们没有特别声明一个流，那么我们的所有操作是在默认的空流中完成的，我们前面的所有例子都是在默认的空流中进行的。空流是没办法管理的。
显式声明的流，我们叫做非空流
非空流可以被显式地创建和管理。如果想要重叠不同的CUDA操作，必须使用非空流。

基于流的异步内核启动和数据传输支持以下类型的粗粒度并发

重叠主机和设备计算
重叠主机计算和主机设备数据传输
重叠主机设备数据传输和设备计算
并发设备计算（多个设备）

CUDA编程有两个“可运算的设备”也就是CPU和GPU这两个东西，这种情况下，他们之间的同步并不是每一步指令都互相通信执行进度的，设备不知道主机在干啥，主机也不是完全知道设备在干啥。数据传输是同步的，也就是主机要等设备接收完数据才干别的。但是内核启动就是异步，无论内核是否完成，主机的应用程序几乎都立即恢复执行。这种内核启动的默认异步行为使它可以直接重叠设备和主机计算。

前面用的cudaMemcpy就是个同步操作，我们还提到过隐式同步——从设备复制结果数据回主机，要等设备执行完。

当然数据传输有异步版本，但是必须显式地设置一个CUDA流来装载它们：

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

特别注意：
执行异步数据传输时，主机端的内存必须是固定的，非分页的。

讲内存模型的时候我们说到过，分配方式：

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

主机虚拟内存中分配的数据在物理内存中是随时可能被移动的，
我们必须确保其在整个生存周期中位置不变，这样在异步操作中才能准确的转移数据，
否则如果操作系统移动了数据的物理地址，那么我们的设备可能还是回到之前的物理地址取数据，
这就会出现未定义的错误。

值得注意的就是cudaMemcpyAsync最后一个参数，stream表示流，一般情况设置为默认流。这个函数和主机是异步的，执行后控制权立刻归还主机，当然我们需要声明一个非空流：

给流命名、声明流的操作应该是：

cudaStream_t a;
cudaStreamCreate(&a);

 定义了一个叫a的流，但是这个流没法用，相当于只有了名字，资源还是要用cudaStreamCreate分配的。

在非空流中执行内核需要在启动核函数的时候加入一个附加的启动配置：
```
kernel_name<<<grid, block, sharedMemSize, stream>>>(argument list);
```
pStream参数就是附加的参数，使用目标流的名字作为参数，比如想把核函数加入到a流中，那么这个stream就变成a。
前面我们为一个流分配资源，当然后面就要回收资源，回收方式：
```
cudaError_t cudaStreamDestroy(cudaStream_t stream);
```
这个回收函数很有意思，由于流和主机端是异步的，你在使用上面指令回收流的资源的时候，很有可能流还在执行，这时候，这条指令会正常执行，但是不会立刻停止流，而是等待流执行完成后，立刻回收该流中的资源。这样做是合理的也是安全的。
当然，我们可以查询流执行的怎么样了，下面两个函数就是帮我们查查我们的流到哪了：
```
cudaError_t cudaStreamSynchronize(cudaStream_t stream);
cudaError_t cudaStreamQuery(cudaStream_t stream);
```
这两条执行的行为非常不同，cudaStreamSynchronize会阻塞主机，直到流完成。cudaStreamQuery则是立即返回，如果查询的流执行完了，那么返回cudaSuccess否则返回cudaErrorNotReady。

下面这段示例代码就是典型多个流中调度CUDA操作的常见模式：

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

第一个for中循环执行了nStreams个流，每个流中都是“复制数据，执行核函数，最后将结果复制回主机”这一系列操作。
下面的图就是一个简单的时间轴示意图，假设nStreams=3，所有传输和核启动都是并发的：

6-1

H2D是主机到设备的内存传输，D2H是设备到主机的内存传输。

显然这些操作没有并发执行，而是错开的，原因是PCIe总线是共享的，当第一个流占据了主线，后来的就一定要等待，等待主线空闲。编程模型和硬件的实际执行时有差距了。

上面同时从主机到设备涉及硬件竞争要等待；具有双工PCIe总线的设备可以重叠两个数据传输，但它们必须在不同的流中以及不同的方向，所以如果是从主机到设备和从设备到主机同时发生，这时候不会产生等待，而是同时进行。

内核并发最大数量也是有极限的，不同计算能力的设备不同，Fermi设备支持16路并发，Kepler支持32路并发。设备上的所有资源都是限制并发数量的原因，比如共享内存，寄存器，本地内存，这些资源都会限制最大并发数。

6.1.2 流调度

从编程模型看，所有流可以同时执行，但是硬件毕竟有限，不可能像理想情况下的所有流都有硬件可以使用，所以硬件上如何调度这些流是我们理解流并发的关键

1. 虚假的依赖关系

在Fermi架构上16路流并发执行，即多达16个网格同时执行，但是所有流最终都是在单一硬件上执行的，Fermi只有一个硬件工作队列，所以他们虽然在编程模型上式并行的，但是在硬件执行过程中是在一个队列中（像串行一样）。

当要执行某个网格的时候CUDA会检测任务依赖关系，如果其依赖于其他结果，那么要等结果出来后才能继续执行。

单一流水线可能会导致虚假依赖关系：

6-2

这个图就是虚假依赖的最准确的描述，我们有三个流，流中的操作相互依赖，比如B要等待A的结果，Z要等待Y的结果，当我们把三个流塞到一个队列中，那么我们就会得到紫色箭头的样子，这个硬件队列中的任务可以并行执行，但是要考虑依赖关系，所以，我们按照顺序会这样执行：

执行A，同时检查B是否有依赖关系，当然此时B依赖于A而A没执行完，所以整个队列阻塞
A执行完成后执行B，同时检查C，发现依赖，等待
B执行完后，执行C同时检查，发现P没有依赖，如果此时硬件有多于资源P开始执行
P执行时检查Q，发现Q依赖P，所以等待

这种一个队列的模式，会产生一种，虽然P依赖B的感觉，虽然不依赖，但是B不执行完，P没办法执行，而所谓并行，只有一个依赖链的头和尾有可能并行，也就是红圈中任务可能并行，而我们的编程模型中设想的并不是这样的。

2. Hyper-Q技术

上面虚假依赖的最好的解决办法就是多个工作队列，这样就从根本上解决了虚假依赖关系，Hyper-Q就是这种技术。

Kepler GPU使用32个硬件工作队列，每个流分配一个工作队列。如果创建的流超过32个，多个流将共享一个硬件工作队列。这样做的结果是可实现全流级并发，并且其具有最小的虚假流间依赖关系。

在这里插入图片描述

6.1.3 流的优先级

3.5以上的设备可以给流优先级.

优先级只影响核函数，不影响数据传输，高优先级的流可以占用低优先级的工作。

下面函数创建一个有指定优先级的流

cudaError_t cudaStreamCreateWithPriority(cudaStream_t* pStream, unsigned int flags,int priority);

这个函数创建了一个具有指定整数优先级的流，并在pStream中返回一个句柄。这个优先级是与pStream中的工作调度相关的。

不同的设备有不同的优先级等级，下面函数可以查询当前设备的优先级分布情况：

cudaError_t cudaDeviceGetStreamPriorityRange(int *leastPriority, int *greatestPriority);

这个函数的返回值存放在leastPriority和greatestPriority中，分别对应于当前设备的最低和最高优先级。

leastPriority表示最低优先级（整数，远离0）
greatestPriority表示最高优先级（整数，数字较接近0）

如果设备不支持优先级返回0

CUDA事件

CUDA事件不同于我们前面介绍的内存事务，不要搞混，事件也是软件层面上的概念。

事件的本质就是一个标记，它与其所在的流内的特定点相关联。

可以使用事件来执行以下两个基本任务：

同步流执行
监控设备的进展

流中的任意点都可以通过API插入事件以及查询事件完成的函数，只有事件所在流中其之前的操作都完成后才能触发事件完成。默认流中设置事件，那么其前面的所有操作都完成时，事件才出发完成。

事件就像一个个路标，其本身不执行什么功能，就像我们最原始测试c语言程序的时候插入的无数多个printf一样。

1. 创建和销毁

事件的声明如下：

cudaEvent_t event;

同样声明完后要分配资源：

cudaError_t cudaEventCreate(cudaEvent_t* event);

回收事件的资源

cudaError_t cudaEventDestroy(cudaEvent_t event);

如果回收指令执行的时候事件还没有完成，那么回收指令立即完成，当事件完成后，资源马上被自动回收。

2. 记录事件和计算运行时间

事件的主要用途就是记录事件之间的时间间隔，检查正在执行的流操作是否已经到达了给定点。

事件通过下面指令添加到CUDA流：

cudaError_t cudaEventRecord(cudaEvent_t event, cudaStream_t stream = 0);

在流中的事件主要作用就是等待前面的操作完成，或者测试指定流中操作完成情况，下面和流类似的事件测试指令（是否出发完成）会阻塞主机线程知道事件被完成。

cudaError_t cudaEventSynchronize(cudaEvent_t event);

同样，也有异步版本，测试一个事件是否可以不用阻塞主机应用程序来完成：

cudaError_t cudaEventQuery(cudaEvent_t event);

这个不会阻塞主机线程，而是直接返回结果和stream版本的类似。

另一个函数用在事件上的是记录两个事件之间的时间间隔：

cudaError_t cudaEventElapsedTime(float* ms, cudaEvent_t start, cudaEvent_t stop);

这个函数记录两个事件start和stop之间的时间间隔，单位毫秒，两个事件不一定是同一个流中。这个时间间隔可能会比实际大一些，因为cudaEventRecord这个函数是异步的，所以加入时间完全不可控，不能保证两个事件之间的间隔刚好是两个事件之间的。

一段简单的记录事件时间间隔的代码

// create two events
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
// record start event on the default stream
cudaEventRecord(start);
// execute kernel
kernel<<<grid, block>>>(arguments);
// record stop event on the default stream
cudaEventRecord(stop);
// wait until the stop event completes
cudaEventSynchronize(stop);
// calculate the elapsed time between two events
float time;
cudaEventElapsedTime(&time, start, stop);
// clean up the two events
cudaEventDestroy(start);
cudaEventDestroy(stop);

这段代码显示，我们的事件被插入到空流中，设置两个事件作为标记，然后记录他们之间的时间间隔。
cudaEventRecord是异步的，所以间隔不准，这是特别要注意的。

6.1.5 流同步

在研究线程并行的时候我们就发现并行这种一旦开始就万马奔腾的模式，想要控制就要让大家到一个固定的位置停下来，就是同步，同步好处是保证代码有可能存在内存竞争的地方降低风险，第二就是相互协调通信，当然坏处就是效率会降低，原因很简单，就是当部分线程等待的时候，设备有一些资源是空闲的，所以这会带来性能损耗。

同样，在流中也有同步，下面我们就研究一下流同步。

从主机的角度，CUDA操作可以分为两类：

内存相关操作：虽然某些内存是同步的，但是他们也有异步版本。
内核启动：内核启动总是异步的

前面我们提到了流的两种类型：

异步流（非空流）：程序员声明的流都是异步流，异步流通常不会阻塞主机
同步流（空流/默认流）：没有显式声明的流式默认同步流，同步流中部分操作会造成阻塞，主机等待，什么都不做，直到某操作完成。

非空流并不都是非阻塞的，其也可以分为两种类型：

阻塞流
非阻塞流

虽然正常来讲，非空流都是异步操作，不存在阻塞主机的情况，但是有时候可能被空流中的操作阻塞。如果一个非空流被声明为非阻塞的，那么没人能阻塞他，如果声明为阻塞流，则会被空流阻塞。

就是非空流有时候可能需要在运行到一半和主机通信，这时候我们更希望他能被阻塞，而不是不受控制，
这样我们就可以自己设定这个流到底受不受控制，也就是是否能被阻塞，下面我们研究如何使用这两种流。

1. 阻塞流和非阻塞流

cudaStreamCreate创建的是阻塞流，意味着里面有些操作会被阻塞，直到空流中默写操作完成。

空流不需要显式声明，而是隐式的，他是阻塞的，跟所有阻塞流同步。

下面这个过程很重要：

当操作A发布到空流中，A执行之前，CUDA会等待A之前的全部操作都发布到阻塞流中，
所有发布到阻塞流中的操作都会挂起，等待，直到在此操作指令之前的操作都完成，才开始执行。

举例：

kernel_1<<<1, 1, 0, stream_1>>>();
kernel_2<<<1, 1>>>();
kernel_3<<<1, 1, 0, stream_2>>>();

上面这段代码，有三个流，两个有名字的，一个空流，我们认为stream_1和stream_2是阻塞流，空流是阻塞的，这三个核函数都在阻塞流上执行，具体过程是

kernel_1被启动，控制权返回主机，
然后启动kernel_2，
但是此时kernel_2 不会并不会马山执行，他会等到kernel_1执行完毕，
同理启动完kernel_2 控制权立刻返回给主机，
主机继续启动kernel_3,这时候kernel_3 也要等待，直到kernel_2执行完，
但是从主机的角度，这三个核都是异步的，启动后控制权马上还给主机。

然后我们就想创建一个非阻塞流，因为我们默认创建的是阻塞版本：

cudaError_t cudaStreamCreateWithFlags(cudaStream_t* pStream, unsigned int flags);

第二个参数就是选择阻塞还是非阻塞版本：

cudaStreamDefault;// 默认阻塞流
cudaStreamNonBlocking: //非阻塞流，对空流的阻塞行为失效。

如果前面的stream_1和stream_2声明为非阻塞的，那么上面的调用方法的结果是三个核函数同时执行。

2. 隐式同步

前面几章核函数计时的时候，我们说过要同步，并且提到过cudaMemcpy 可以隐式同步，也介绍了

cudaDeviceSynchronize;
cudaStreamSynchronize;
cudaEventSynchronize;

这几个也是同步指令，可以用来同步不同的对象，这些是显式的调用的；与上面的隐式不同。

隐式同步的指令其最原始的函数功能并不是同步，所以同步效果是隐式的。
这个我们需要非常注意，忽略隐式同步会造成性能下降。

所谓同步就是阻塞的意思，被忽视的隐式同步就是被忽略的阻塞，隐式操作常出现在内存操作上，比如：

锁页主机内存分布
设备内存分配
设备内存初始化
同一设备两地址之间的内存复制
一级缓存，共享内存配置修改

这些操作都要时刻小心，因为他们带来的阻塞非常不容易察觉

3. 显式同步

显式同步相比就更加光明磊落了，因为一条指令就一个作用，没啥副作用，
CUDA运行时在网格级支持显式同步CUDA程序的几种方法：

同步设备
同步流
同步流中的事件
使用事件跨流同步

下面的函数就可以阻塞主机线程，直到设备完成所有操作，尽量少用，这个会拖慢效率。

cudaError_t cudaDeviceSynchronize(void);

然后是流版本的，我们可以同步流，使用下面两个函数：

cudaError_t cudaStreamSynchronize(cudaStream_t stream);
cudaError_t cudaStreamQuery(cudaStream_t stream);

这两个函数，第一个是同步流的，阻塞主机直到完成，
第二个可以完成非阻塞流测试。也就是测试一下这个流是否完成。

事件的作用就是在流中设定一些标记用来同步，和检查是否执行到关键点位（事件位置），也是用类似的函数

cudaError_t cudaEventSynchronize(cudaEvent_t event);
cudaError_t cudaEventQuery(cudaEvent_t event);

这两个函数的性质和上面的非常类似。

事件提供了一个跨流同步的方法：

cudaError_t cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event);

这条命令的含义是，指定的流要等待指定的事件，事件完成后流才能继续，
这个事件可以在这个流中，也可以不在，当在不同的流的时候，这个就是实现了跨流同步。

在这里插入图片描述

4. 可配置事件

CDUA提供了一种控制事件行为和性能的函数：

cudaError_t cudaEventCreateWithFlags(cudaEvent_t* event, unsigned int flags);

其中参数是：

cudaEventDefault
cudaEventBlockingSync
cudaEventDisableTiming
cudaEventInterprocess

其中cudaEventBlockingSync指定使用cudaEventSynchronize同步会造成阻塞调用线程。

cudaEventSynchronize默认是使用cpu周期不断重复查询事件状态，而当指定了事件是cudaEventBlockingSync的时候，会将查询放在另一个线程中，而原始线程继续执行，直到事件满足条件，才会通知原始线程，这样可以减少CPU的浪费；但是由于通讯的时间，会造成一定的延迟。
cudaEventDisableTiming表示事件不用于计时，可以减少系统不必要的开支也能提升cudaStreamWaitEvent和cudaEventQuery的效率
cudaEventInterprocess表明可能被用于进程之间的事件

6.2 并发内核执行

介绍了流，事件和同步等的概念，以及一些函数的用法，接下来的几个例子，介绍并发内核的几个基本问题，包括不限于以下几个方面：

使用深度优先或者广度优先方法的调度工作
调整硬件工作队列
在Kepler设备和Fermi设备上避免虚假的依赖关系
检查默认流的阻塞行为
在非默认流之间添加依赖关系
检查资源使用是如何影响并发的

6.2.1 非空流中的并发内核

本文我们开始使用NVIDIA提供的另一个可视化工具nvvp进行性能分析，其最大用途在于可视化并发核函数的执行，

第一个例子中我们就能清楚地看到各个核函数是如何执行的，本例子中使用了同一个核函数，并将其复制多份，并确保每个核函数的计算要消耗足够的时间，保证执行过程能够被性能分析工具准确的捕捉到。
我们的核函数是：

__global__ void kernel_1()
{
    double sum=0.0;
    for(int i=0;i<N;i++)
        sum=sum+tan(0.1)*tan(0.1);
}
__global__ void kernel_2()
{
    double sum=0.0;
    for(int i=0;i<N;i++)
        sum=sum+tan(0.1)*tan(0.1);
}
__global__ void kernel_3()
{
    double sum=0.0;
    for(int i=0;i<N;i++)
        sum=sum+tan(0.1)*tan(0.1);
}
__global__ void kernel_4()
{
    double sum=0.0;
    for(int i=0;i<N;i++)
        sum=sum+tan(0.1)*tan(0.1);
}

接着我们按照上节课的套路，创建流，把不同的核函数或者指令放到不同的流中，然后看一下他们的表现。

我们本章主要关注主机代码，下面是创建流的代码：

cudaStream_t *stream=(cudaStream_t*)malloc(n_stream*sizeof(cudaStream_t));
for(int i=0;i<n_stream;i++)
{
    cudaStreamCreate(&stream[i]);
}

首先声明一个流的头结构，是malloc的注意后面要free掉
然后为每个流的头结构分配资源，也就是Create的过程，这样我们就有n_stream个流可以使用了.

接着，我们添加核函数到流，并观察运行效果

dim3 block(1);
dim3 grid(1);
cudaEvent_t start,stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start);
for(int i=0;i<n_stream;i++)
{
    kernel_1<<<grid,block,0,stream[i]>>>();
    kernel_2<<<grid,block,0,stream[i]>>>();
    kernel_3<<<grid,block,0,stream[i]>>>();
    kernel_4<<<grid,block,0,stream[i]>>>();
}
cudaEventRecord(stop);
CHECK(cudaEventSynchronize(stop));
float elapsed_time;
cudaEventElapsedTime(&elapsed_time,start,stop);
printf("elapsed time:%f ms\n",elapsed_time);

这个循环是将每个核函数都放入不同的流之中，也就是假设我们有10个流，那么这10个流中每个流都要按照上面的顺序执行这4个核函数。

注意如果没有cudaEventSynchronize(stop)，nvvp将会无法运行，因为所有这些都是异步操作，不会等到操作完再返回，而是启动后自动把控制权返回主机，如果没有一个阻塞指令，主机进程就会执行完毕推出，这样就跟设备失联了，nvvp也会相应的报错。

然后我们创建两个事件，然后记录事件之间的时间间隔。这个间隔是不太准确的，因为是异步的。

使用nvvp检测，

$ nvvp ./main

结果如下：

在这里插入图片描述

6.2.2 Fermi GPU 上的虚假依赖关系

虚假依赖我们在上文中讲到过了，这种情况通常出现在只有在比较古老的Fermi架构上出现，原因是其只有一个硬件工作队列，由于我们现在很难找到Fermi架构的GPU了，所以，只能看看书上给出的nvvp结果图了：

nvvp-2

如果你手头只有老机器，这种虚假依赖关系也是可以解决的，原理就是使用广度优先的方法，组织各任务的方式如下：

// dispatch job with breadth first way
for (int i = 0; i < n_streams; i++)
kernel_1<<<grid, block, 0, streams[i]>>>();
for (int i = 0; i < n_streams; i++)
kernel_2<<<grid, block, 0, streams[i]>>>();
for (int i = 0; i < n_streams; i++)
kernel_3<<<grid, block, 0, streams[i]>>>();
for (int i = 0; i < n_streams; i++)
kernel_4<<<grid, block, 0, streams[i]>>>();

采用广度优先顺序可以确保工作队列中相邻的任务来自于不同的流。因此，任何相邻的任务对之间都不会再有虚假的依赖关系，从而得以实现并发内核执行。

逻辑图:
在这里插入图片描述

广度优先的nvvp结果是：

在这里插入图片描述

6.2.3 使用OpenMP的调度操作

前面的示例中，使用单一的主机线程将异步CUDA操作调度到多个流中。本节的示例将使用多个主机线程将操作调度到多个流中，并使用一个线程来管理每一个流。

OpenMP是一种非常好用的并行工具，比pthread更加好用，但是没有pthread那么灵活，这里我们不光要让核函数或者设备操作用多个流处理，同时也让主机在多线程下工作，我们尝试使用每个线程来操作一个流：

omp_set_num_thread(n_stream);
#pragma omp parallel
    {
        int i=omp_get_thread_num();
        kernel_1<<<grid,block,0,stream[i]>>>();
        kernel_2<<<grid,block,0,stream[i]>>>();
        kernel_3<<<grid,block,0,stream[i]>>>();
        kernel_4<<<grid,block,0,stream[i]>>>();
    }

解释下代码

omp_set_num_thread(n_stream);
#pragma omp parallel

调用OpenMP的API创建n_stream个线程，然后宏指令告诉编译器下面大括号中的部分就是每个线程都要执行的部分，有点类似于核函数，或者叫做并行单元。
omp_get_thread_num函数为每个主机线程返回唯一一个线程ID，将该ID用作streams数组中
的索引，用来创建OpenMP线程和CUDA流间的一对一映射。

nvcc  -O3 -Xcompiler -fopenmp stream_omp.cu -o stream_omp -lgomp

6.2.4 用环境变量调整流行为

Kepler支持的最大Hyper-Q 工作队列数是32 ，但是在默认情况下并不是全部开启，而是被限制成8个，原因是每个工作队列只要开启就会有资源消耗，如果用不到32个可以把资源留给需要的8个队列，修改这个配置的方法是修改主机系统的环境变量。

对于Linux系统中，修改方式如下：

#For Bash or Bourne Shell:
export CUDA_DEVICE_MAX_CONNECTIONS=32
#For C-Shell:
setenv CUDA_DEVICE_MAX_CONNECTIONS 32

另一种修改方法是直接在程序里写，这种方法更好用通过底层驱动修改硬件配置：

setenv("CUDA_DEVICE_MAX_CONNECTIONS", "32", 1);

每个CUDA流都会被映射到单一的CUDA设备连接中。如果流的数量超过了硬件连接的数量，多个流将共享一个连接。当多个流共享相同的硬件工作队列时，可能会产生虚假的依赖关系。

然后我们把前面的深度优先的代码改一下，加入上面这句指令,并把n_stream改成16个流，8个工作队列的结果：
在这里插入图片描述
可以用广度优先顺序调度内核去除了虚假的依赖关系。

6.2.5 GPU资源的并发限制

限制内核并发数量的最根本的还是GPU上面的资源，资源才是性能的极限，性能最高无非是在不考虑算法进化的前提下，资源利用率最高的结果。

当每个内核的线程数增加的时候，内核级别的并行数量就会下降，比如，我们把

dim3 block(1);
dim3 grid(1);

升级到

dim3 block(16,32);
dim3 grid(32);

4个流，nvvp结果是：
在这里插入图片描述

6.2.6 默认流的阻塞行为

默认流也就是空流，对于非空流中的阻塞流是有阻塞作用的，

首先我们没有声明流的那些GPU操作指令，核函数是在空流上执行的，空流是阻塞流，
同时我们声明定义的流如果没有特别指出，声明的也是阻塞流，换句话说，这些流的共同特点，无论空流与非空流，都是阻塞的。
那么这时候空流（默认流）对非空流的阻塞操作就要注意一下了。

for(int i=0;i<n_stream;i++)
{
    kernel_1<<<grid,block,0,stream[i]>>>();
    kernel_2<<<grid,block,0,stream[i]>>>();
    kernel_3<<<grid,block>>>();
    kernel_4<<<grid,block,0,stream[i]>>>();
}

注意，kernel_3是在空流（默认流）上的，从NVVP的结果中可以看出，所有kernel_3 启动以后，所有其他的流中的操作全部被阻塞。

6.2.7 创建流间依赖关系

流之间的虚假依赖关系是需要避免的，而经过我们设计的依赖又可以保证流之间的同步性，避免内存竞争，这时候我们要使用的就是事件这个工具了，

我们可以让某个特定流等待某个特定的事件，这个事件可以在任何流中，只有此事件完成才能进一步执行等待此事件的流继续执行。
这种事件往往不用于计时，所以可以在生命的时候声明成 cudaEventDisableTiming 的同步事件：

cudaEvent_t * event=(cudaEvent_t *)malloc(n_stream*sizeof(cudaEvent_t));
for(int i=0;i<n_stream;i++)
{
    cudaEventCreateWithFlag(&event[i],cudaEventDisableTiming);
}

在流中加入指令:

for(int i=0;i<n_stream;i++)
{
    kernel_1<<<grid,block,0,stream[i]>>>();
    kernel_2<<<grid,block,0,stream[i]>>>();
    kernel_3<<<grid,block,0,stream[i]>>>();
    kernel_4<<<grid,block,0,stream[i]>>>();
    cudaEventRecord(event[i],stream[i]);
    cudaStreamWaitEvent(stream[n_stream-1],event[i],0);
}

这时候，最后一个流（streams[n_streams-1]）都会等到前面所有流中的事件完成，自己才会完成

在这里插入图片描述

我想静静，

关注

1
点赞
踩
9

收藏

觉得还不错? 一键收藏
0
评论
CUDA学习之流和并发--part1

文章目录6.1 流和事件概述6.1.1 CUDA流6.1.2 流调度1. 虚假的依赖关系2. Hyper-Q技术6.1.3 流的优先级CUDA事件1. 创建和销毁2. 记录事件和计算运行时间6.1.5 流同步1. 阻塞流和非阻塞流2. 隐式同步3. 显式同步4. 可配置事件6.2 并发内核执行6.2.1 非空流中的并发内核6.2.2 Fermi GPU 上的虚假依赖关系6.2.3 使用OpenMP的调度操作6.2.4 用环境变量调整流行为6.2.5 GPU资源的并发限制6.2.6 默认流的阻塞行为6.2.7
复制链接

扫一扫