title: 【CUDA 基础】6.1 流和事件概述
categories:
- CUDA
- Freshman
tags:
- 流
- 事件
toc: true
date: 2018-06-10 21:45:15
Abstract: 本文介绍CUDA中流和事件的理论描述。
Keywords: 流，事件
开篇废话
今天没废话，废话太多影响效率?
前面几章我们一直围绕GPU设备展开，我们的代码除了在核函数的配置的部分研究过主机端执行的代码，其他部分基本都是在设备代码上进行的，这一章我们就从主机端来讲讲如何优化CUDA应用。
CUDA流：一系列异步CUDA操作，比如我们常见的套路，在主机端分配设备主存（cudaMalloc），主机向设备传输数据（cudaMemcpy），核函数启动，复制数据回主机（Memcpy）这些操作中有些是异步的，执行顺序也是按照主机代码中的顺序执行的（但是异步操作的结束不一定是按照代码中的顺序的）。
流能封装这些异步操作，并保持操作顺序，允许操作在流中排队。保证其在前面所有操作启动之后启动，有了流，我们就能查询排队状态了。
我们上面举得一般情况下的操作基本可以分为以下三种：
- 主机与设备间的数据传输
- 核函数启动
- 其他的由主机发出的设备执行的命令
流中的操作相对于主机来说总是异步的，CUDA运行时决定何时可以在设备上执行操作。我们要做的就是控制这些操作在其结果出来之前，不启动需要调用这个结果的操作。
一个流中的不同操作有着严格的顺序。但是不同流之间是没有任何限制的。多个流同时启动多个内核，就形成了网格级别的并行。
CUDA流中排队的操作和主机都是异步的，所以排队的过程中并不耽误主机运行其他指令，所以这就隐藏了执行这些操作的开销。
CUDA编程的一个典型模式是，也就是我们上面讲到的一般套路：
- 将输入数据从主机复制到设备上
- 在设备上执行一个内核
- 将结果从设备移回主机
一般的生产情况下，内核执行的时间要长于数据传输，所以我们前面的例子大多是数据传输更耗时，这是不实际的。当重叠核函数执行和数据传输操作，可以屏蔽数据移动造成的时间消耗，当然正在执行的内核的数据需要提前复制到设备上的，这里说的数据传输和内核执行是同时操作的是指当前传输的数据是接下来流中的内核需要的。这样总的执行时间就被缩减了。
流在CUDA的API调用可以实现流水线和双缓冲技术。
CUDA的API也分为同步和异步的两种：
- 同步行为的函数会阻塞主机端线程直到其完成
- 异步行为的函数在调用后会立刻把控制权返还给主机。
异步行为和流式构建网格级并行的支柱。
虽然我们从软件模型上提出了流，网格级并行的概念，但是说来说去我们能用的就那么一个设备，如果设备空闲当然可以同时执行多个核，但是如果设备已经跑满了，那么我们认为并行的指令也必须排队等待——PCIe总线和SM数量是有限的，当他们被完全占用，流是没办法做什么的，除了等待
我们接下来就要研究多种计算能力的设备上的流是如何运行的。
CUDA流
我们的所有CUDA操作都是在流中进行的，虽然我们可能没发现，但是有我们前面的例子中的指令，内核启动，都是在CUDA流中进行的，只是这种操作是隐式的，所以肯定还有显式的，所以，流分为：
- 隐式声明的流，我们叫做空流
- 显式声明的流，我们叫做非空流
如果我们没有特别声明一个流，那么我们的所有操作是在默认的空流中完成的，我们前面的所有例子都是在默认的空流中进行的。
空流是没办法管理的，因为他连个名字都没有，似乎也没有默认名，所以当我们想控制流，非空流是非常必要的。
基于流的异步内核启动和数据传输支持以下类型的粗粒度并发
- 重叠主机和设备计算
- 重叠主机计算和主机设备数据传输
- 重叠主机设备数据传输和设备计算
- 并发设备计算（多个设备）
CUDA编程和普通的C++不同的就是，我们有两个“可运算的设备”也就是CPU和GPU这两个东西，这种情况下，他们之间的同步并不是每一步指令都互相通信执行进度的，设备不知道主机在干啥，主机也不是完全知道设备在干啥。但是数据传输是同步的，也就是主机要等设备接收完数据才干别的，也就是说你爸给你寄了一袋大米，然后老人家啥也不做，拨通电话跟你保持通话不停的问你收到了么？直到你回答收到了，这就是同步的。内核启动就是异步的，你爸爸又要给你钱花，去银行给你汇了五百块钱，银行说第二天到账，他就可以回家该干嘛干嘛了，而不需要在银行等一晚，第二天你收到了，打个电话说一声就行了，这就是异步的。异步操作，可以重叠主机计算和设备计算。
前面用的cudaMemcpy就是个同步操作，我们还提到过隐式同步——从设备复制结果数据回主机，要等设备执行完。当然数据传输有异步版本：
cudaError_t cudaMemcpyAsync(void* dst, const void* src, size_t count,cudaMemcpyKind kind, cudaStream_t stream = 0);
值得注意的就是最后一个参数，stream表示流，一般情况设置为默认流，这个函数和主机是异步的，执行后控制权立刻归还主机，当然我们需要声明一个非空流：
cudaError_t cudaStreamCreate(cudaStream_t* pStream);
这样我们就有一个可以被管理的流了，这段代码是创建了一个流，有C++经验的人能看出来，这个是为一个流分配必要资源的函数，给流命名声明流的操作应该是：
cudaStream_t a;
定义了一个叫a的流，但是这个流没法用，相当于只有了名字，资源还是要用cudaStreamCreate分配的。
接下来必须要特别注意：
执行异步数据传输时，主机端的内存必须是固定的，非分页的！！
执行异步数据传输时，主机端的内存必须是固定的，非分页的！！
执行异步数据传输时，主机端的内存必须是固定的，非分页的！！
讲内存模型的时候我们说到过，分配方式：
cudaError_t cudaMallocHost(void **ptr, size_t size); cudaError_t cudaHostAlloc(void **pHost, size_t size, unsigned int flags);
主机虚拟内存中分配的数据在物理内存中是随时可能被移动的，我们必须确保其在整个生存周期中位置不变，这样在异步操作中才能准确的转移数据，否则如果操作系统移动了数据的物理地址，那么我们的设备可能还是回到之前的物理地址取数据，这就会出现未定义的错误。
在非空流中执行内核需要在启动核函数的时候加入一个附加的启动配置：
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，所有传输和核启动都是并发的：
H2D是主机到设备的内存传输，D2H是设备到主机的内存传输。显然这些操作没有并发执行，而是错开的，原因是PCIe总线是共享的，当第一个流占据了主线，后来的就一定要等待，等待主线空闲。编程模型和硬件的实际执行时有差距了。
上面同时从主机到设备涉及硬件竞争要等待，如果是从主机到设备和从设备到主机同时发生，这时候不会产生等待，而是同时进行。
内核并发最大数量也是有极限的，不同计算能力的设备不同，Fermi设备支持16路并发，Kepler支持32路并发。设备上的所有资源都是限制并发数量的原因，比如共享内存，寄存器，本地内存，这些资源都会限制最大并发数。
