CUDA ---- Stream and Event

时间 2019-11-20

标签 cuda stream event 栏目 Apache 繁體版

原文原文链接

Streamhtml

通常来讲，cuda c并行性表如今下面两个层面上：编程

Kernel level
Grid level

到目前为止，咱们讨论的一直是kernel level的，也就是一个kernel或者一个task由许多thread并行的执行在GPU上。Stream的概念是相对于后者来讲的，Grid level是指多个kernel在一个device上同时执行。api

Stream和event简介

Cuda stream是指一堆异步的cuda操做，他们按照host代码调用的顺序执行在device上。Stream维护了这些操做的顺序，并在全部预处理完成后容许这些操做进入工做队列，同时也能够对这些操做进行一些查询操做。这些操做包括host到device的数据传输，launch kernel以及其余的host发起由device执行的动做。这些操做的执行老是异步的，cuda runtime会决定这些操做合适的执行时机。咱们则可使用相应的cuda api来保证所取得结果是在全部操做完成后得到的。同一个stream里的操做有严格的执行顺序，不一样的stream则没有此限制。并发

因为不一样stream的操做是异步执行的，就能够利用相互之间的协调来充分发挥资源的利用率。典型的cuda编程模式咱们已经熟知了：异步

将输入数据从host转移到device
在device上执行kernel
将结果从device上转移回host

在许多状况下，花费在执行kernel上的时间要比传输数据多得多，因此很容易想到将cpu和gpu之间的沟通时间隐藏在其余kernel执行过程当中，咱们能够将数据传输和kernel执行放在不一样的stream中来实现此功能。Stream能够用来实现pipeline和双buffer（front-back）渲染。async

Cuda API可分为同步和异步两类，同步函数会阻塞host端的线程执行，异步函数会马上将控制权返还给host从而继续执行以后的动做。异步函数和stream是grid level并行的两个基石。函数

从软件角度来看，不一样stream中的不一样操做能够并行执行，可是硬件角度却不必定如此。这依赖于PCIe连接或者每一个SM可得到的资源，不一样的stream仍然须要等待别的stream来完成执行。下面会简单介绍在不一样CC版本下，stream在device上的行为。性能

Cuda Streams

全部的cuda操做（包括kernel执行和数据传输）都显式或隐式的运行在stream中，stream也就两种类型，分别是：测试

隐式声明stream（NULL stream）
显示声明stream（non-NULL stream）

默认状况下是NULL stream，在以前未涉及到stream的博文中，都是该类型。若是显式的声明一个stream就是non-NULL stream了。优化

异步且基于stream的kernel执行和数据传输可以实现如下几种类型的并行：

Host运算操做和device运算操做并行
Host运算操做和host到device的数据传输并行
Host到device的数据传输和device运算操做并行
Device内的运算并行

下面代码是以前常见的使用形式，默认使用NULL stream:

cudaMemcpy(..., cudaMemcpyHostToDevice);
kernel<<<grid, block>>>(...);
cudaMemcpy(..., cudaMemcpyDeviceToHost);

从device角度看，全部者三个操做都是使用的默认stream，而且按照代码从上到下的顺序依次执行，device自己是不知道其余的host操做怎样执行的。从host角度来看，数据传输都是同步的而且会一直等待，直到操做完成。不过不一样于数据传输，Kernel的launch是异步的，host差很少马上就能从新获得控制权，不用管kernel是否执行完毕，从而进行下一步动做。很明显，这种异步行为有助于重叠device和host之间的运算时间。

上文内容在以前博文都有涉及，这里特别说明的是数据传输，它也是能够异步执行的，这就用到了本次讲的stream，咱们必须显示的声明一个stream来分派它的执行。下面版本是异步版本的cudaMemcpy：

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

注意新增长的最后一个参数。这样，在host issue了这个函数给device执行后，控制权能够马上返还给host。上面代码使用了默认stream，若是要声明一个新的stream则使用下面的API定义一个：

cudaError_t cudaStreamCreate(cudaStream_t* pStream);

这样就定义了一个可使用在cuda异步API函数中stream。使用该函数的一个比较常见的错误，或者说容易引发混乱的地方是，这个函数返回的error code多是上一次调用异步函数产生的。也就是说，函数返回error并非调用该函数产生error的必要条件。

当执行一次异步数据传输时，咱们必须使用pinned（或者non-pageable）memory。Pinned memory的分配以下，具体请参见前面博文：

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

经过在将该内存pin到host的虚拟内存上，就能够将该memory的物理位置强制分配到CPU内存中以便使之在整个程序生命周期中保持不变。不然的话，操做系统可能会在任意时刻改变该host端的虚拟内存对应的物理地址。假设异步数据传输函数没有使用pinned host memory的话，操做系统就可能将数据从一块物理空间移动到另外一块物理空间（由于是异步的，CPU在执行其余的动做就可能影响这块数据），而此时cuda runtime正在执行数据的传输，这会致使不肯定的行为。

在执行kernel时要想设置stream的话，也是很简单的，一样只要加一个stream参数就好：

kernel_name<<<grid, block, sharedMemSize, stream>>>(argument list);

// 非默认的stream声明
cudaStream_t stream;
// 初始化
cudaStreamCreate(&stream);
// 资源释放
cudaError_t cudaStreamDestroy(cudaStream_t stream);

当执行资源释放的时候，若是仍然有stream的工做没干完，那么虽然该函数仍然会马上返回，可是相关的工做作完后，这些资源才会自动的释放掉。

因为全部stram的执行都是异步的，就须要一些API在必要的时候作同步操做：

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

第一个会强制host阻塞等待，直至stream中全部操做完成为止；第二个会检查stream中的操做是否所有完成，即便有操做没完成也不会阻塞host。若是全部操做都完成了，则返回cudaSuccess，不然返回cudaErrorNotReady。

下面看一下一个代码片断来帮助理解：

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

该段代码使用了三个stream，数据传输和kernel运算都被分配在了这几个并发的stream中。

上图就跟流水线同样差很少的道理，很少说。须要注意的是，上图中数据传输的操做并非并行执行的，即便他们是在不一样的stream中。按惯例，这种状况确定就是硬件资源的锅了，硬件资源就那么些，软件层面作的优化无非就是尽可能让全部硬件资源一刻不停的被利用起来（万恶的资本主义，嗯……），而这里就是PCIe卡了瓶颈。固然从编程角度来看，这些操做依然是相互独立的，只是他们要共享硬件资源，就不得不是串行的。有两个PCIe就能够重叠这两次数据传输操做，不过也是要保证不一样的stream和不一样的传输方向。

最大并发kernel数目是依赖于device自己的，Fermi支持16路并行，Kepler是32。并行数是受限于shared memory，寄存器等device资源。

Stream Scheduling

概念上来讲，全部stream是同时运行的。可是，事实上一般并不是如此。

False Dependencies

尽管Fermi最高支持16路并行，可是在物理上，全部stream是被塞进硬件上惟一一个工做队列来调度的，当选中一个grid来执行时，runtime会查看task的依赖关系，若是当前task依赖前面的task，该task就会阻塞，因为只有一个队列，后面的都会跟着等待，即便后面的task是别的stream上的任务。就以下图所示：

C和P以及R和X是能够并行的，由于他们在不一样的stream中，可是ABC，PQR以及XYZ却不行，好比，在B没完成以前，C和P都在等待。

Hyper-Q

伪依赖的状况在Kepler系列里获得了解决，采用的一种叫Hyper-Q的技术，简单粗暴的理解就是，既然工做队列不够用，那就增长好了，因而Kepler上出现了32个工做队列。该技术也实现了TPC上能够同时运行compute和graphic的应用。固然，若是超过32个stream被建立了，依然会出现伪依赖的状况。

Stream Priorities

对于CC3.5及以上版本，stream能够有优先级的属性：

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

该函数建立一个stream，赋予priority的优先级，高优先级的grid能够抢占低优先级执行。不过优先级属性只对kernel有效，对数据传输无效。此外，若是设置的优先级超出了可设置范围，则会自动设置成最高或者最低。有效可设置范围可用下列函数查询：

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

顾名思义，leastPriority是下限，gretestPriority是上限。老规矩，数值较小则拥有较高优先级。若是device不支持优先级设置，则这两个值都返回0。

Cuda Events

Event是stream相关的一个重要概念，其用来标记strean执行过程的某个特定的点。其主要用途是：

同步stream执行
操控device运行步调

Cuda api提供了相关函数来插入event到stream中和查询该event是否完成（或者叫知足条件？）。只有当该event标记的stream位置的全部操做都被执行完毕，该event才算完成。关联到默认stream上的event则对全部的stream有效。

Creation and Destruction

// 声明
cudaEvent_t event;
// 建立
cudaError_t cudaEventCreate(cudaEvent_t* event);
// 销毁
cudaError_t cudaEventDestroy(cudaEvent_t event);

同理streeam的释放，在调用该函数的时候，若是相关操做没完成，则会在操做完成后自动释放资源。

Recording Events and Mesuring Elapsed Time

Events标记了stream执行过程当中的一个点，咱们就能够检查正在执行的stream中的操做是否到达该点，咱们能够把event当成一个操做插入到stream中的众多操做中，当执行到该操做时，所作工做就是设置CPU的一个flag来标记表示完成。下面函数将event关联到指定stream。

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

等待event会阻塞调用host线程，同步操做调用下面的函数：

cudaError_t cudaEventSynchronize(cudaEvent_t event);

该函数相似于cudaStreamSynchronize，只不过是等待一个event而不是整个stream执行完毕。咱们同时可使用下面的API来测试event是否完成，该函数不会阻塞host：

cudaError_t cudaEventQuery(cudaEvent_t event);

该函数相似cudaStreamQuery。此外，还有专门的API能够度量两个event之间的时间间隔：

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

返回start和stop之间的时间间隔，单位是毫秒。Start和stop没必要关联到同一个stream上，可是要注意，若是两者任意一个关联到了non-NULL stream上，时间间隔可能要比指望的大。这是由于cudaEventRecord是异步发生的，咱们没办法保证度量出来的时间刚好就是两个event之间，因此只是想要gpu工做的时间间隔，则stop和strat都关联到默认stream就行了。

下面代码简单展现了如何使用event来度量时间：

// 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);

Stream Synchronization

因为全部non-default stream的操做对于host来讲都是非阻塞的，就须要相应的同步操做。

从host的角度来看，cuda操做能够被分为两类：

Memory相关的操做
Kernel launch

Kernel launch对于host来讲都是异步的，许多memory操做则是同步的，好比cudaMemcpy，可是，cuda runtime也会提供异步函数来执行memory操做。

咱们已经知道Stream能够被分为同步（NULL stream）和异步（non-NULL stream）两种，同步异步是针对host来说的，异步stream不会阻塞host的执行，而大多数同步stream则会阻塞host，不过kernel launch例外，不会阻塞host。

此外，异步stream又能够被分为阻塞和非阻塞两种，阻塞非阻塞是异步stream针对同步stream来说的。异步stream若是是阻塞stream，那么同步stream会阻塞该异步stream中的操做。若是异步stream是非阻塞stream，那么该stream不会阻塞同步stream中的操做（有点绕……）。

阻塞和非阻塞stream

使用cudaStreamCreate建立的是阻塞stream，也就是说，该stream中执行的操做会被早先执行的同步stream阻塞。一般来讲，当issue一个NULL stream时，cuda context会等待以前全部阻塞stream完成后才执行该NULL stream，固然全部阻塞stream也会等待以前的NULL stream完成才开始执行。

例如：

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

从device角度来讲，这三个kernel是串行依次执行的，固然从host角度来讲，倒是并行非阻塞的。除了经过cudaStreamCreate生成的阻塞stream外，咱们还能够经过下面的API配置生成非阻塞stream：

cudaError_t cudaStreamCreateWithFlags(cudaStream_t* pStream, unsigned int flags);
// flag为如下两种，默认为第一种，非阻塞即是第二种。
cudaStreamDefault: default stream creation flag (blocking)
cudaStreamNonBlocking: asynchronous stream creation flag (non-blocking)

若是以前的kernel_1和kernel_3的stream被定义成第二种，就不会被阻塞。

Implicit Synchronization

Cuda有两种类型的host和device之间同步：显式和隐式。咱们以前已经了解到显式同步API有：

cudaDeviceSynchronize
cudaStreamSynchronize
cudaEventSynchronize

这三个函数由host显式的调用，在device上执行。

隐式同步咱们也了解过，好比cudaMemcpy就会隐式的同步device和host，由于该函数同步做用只是数据传输的反作用，因此称为隐式。了解这些隐式同步是很中要的，由于不经意的调用这样一个函数可能会致使性能急剧下降。

隐式同步是cuda编程中比较特殊状况，由于隐式同步行为可能会致使意外的阻塞行为，一般发生在device端。许多memory相关的操做都会影响当前device的操做，好比：

A page-locked host memory allocation
A device memory allocation
A device memset
A memory copy between two addresses on the same device
A modification to the L1/shared memory confi guration

Explicit Synchronization

从grid level来看显式同步方式，有以下几种：

Synchronizing the device
Synchronizing a stream
Synchronizing an event in a stream
Synchronizing across streams using an event

咱们可使用以前提到过的cudaDeviceSynchronize来同步该device上的全部操做。该函数会致使host等待全部device上的运算或者数据传输操做完成。显而易见，该函数是个heavyweight的函数，咱们应该尽可能减小这类函数的使用。

经过使用cudaStreamSynchronize可使host等待特定stream中的操做所有完成或者使用非阻塞版本的cudaStreamQuery来测试是否完成。

Cuda event能够用来实现更细粒度的阻塞和同步，相关函数为cudaEventSynchronize和cudaEventSynchronize，用法相似stream相关的函数。此外，cudaStreamWaitEvent提供了一种灵活的方式来引入stream之间的依赖关系：

cudaError_t cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event);

该函数会指定该stream等待特定的event，该event能够关联到相同或者不一样的stream，对于不一样stream的状况，以下图所示：

Stream2会等待stream1中的event完成后继续执行。

Configurable Events

Event的配置可用下面函数：

cudaError_t cudaEventCreateWithFlags(cudaEvent_t* event, unsigned int flags);
cudaEventDefault
cudaEventBlockingSync
cudaEventDisableTiming
cudaEventInterprocess

cudaEventBlockingSync说明该event会阻塞host。cudaEventSynchronize默认行为是使用CPU时钟来固定的查询event状态。使用cudaEventBlockingSync，调用线程会进入休眠，将控制权交给其余线程或者进程，直到event完成为止。可是这样会致使少许的CPU时钟浪费，也会增长event完成和唤醒线程的之间的时间消耗。

cudaEventDisableTiming指定event只能用来同步，而且不须要记录计时数据。这样扔掉记录时间戳的消耗能够提升cuudaStreamWaitEvent和cudaEventQuery的调用性能。

cudaEventInterprocess指定event能够被用来做为inter-process event。

NVIDIA CUDA板块：https://developer.nvidia.com/cuda-zone

CUDA在线文档：http://docs.nvidia.com/cuda/index.html#

转载原文注明：http://www.cnblogs.com/1024incn/p/5891051.html

填坑中~~