【CUDA】Stream and Event

一、stream

1.1 认识stream

1.2 stream调度

1.2.1 伪依赖

1.2.2 Hyper-Q

1.2.3 stream优先级

1.3 使用示例

1.4 流同步

1.4.1 阻塞与非阻塞stream

1.4.2 隐式同步

1.4.3 显示同步

二、Event

2.1 创造与销毁

2.2 记录事件和测量经过的时间

2.3 可配置Events

一、stream

一般而言，cuda并行性表现在下面两个层面上：

Kernel level
Grid level

kernel level，即一个 kernel 或者一个 task 由许多 thread 并行的执行在GPU上。Stream的概念是相对于后者来说的，Grid level是指多个 kernel 在一个 device 上同时执行（一个 Grid 中的 Block 可以在多个 SM 中执行）

1.1 认识stream

流是指一系列指令，且 CUDA 具有默认流。默认情况下，CUDA 核函数会在默认流中运行

在一个流中排队的所有命令都必须在该流中的下一个命令开始执行之前完成（或者至少达到一个可以安全执行下一个命令的状态）。CUDA流提供了一种方式来组织命令的执行顺序，确保一个流中的命令按顺序执行，而不同的流可以并发执行，但每个流内部保持顺序性

不同的非默认流中的核函数可并发执行。默认流较为特殊，其执行任何操作期间，任何非默认流中皆不可同时执行任何操作，默认流将等待非默认流全部执行完毕后再开始运行，在其执行完毕后，其他非默认流才能开始执行

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

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

下面代码是之前常见的使用形式，使用默认 stream：

cudaMemcpy(..., cudaMemcpyHostToDevice);
kernel<<<grid, block>>>(...);
cudaDeviceSynchronize();
// ... ... host计算
cudaMemcpy(..., cudaMemcpyDeviceToHost);

从 device 角度看，所有者三个操作都是使用的默认stream，并且按照代码从上到下的顺序依次执行，device 本身是不知道其他的 host 操作怎样执行的

从 host 角度来看，数据传输都是同步的并且会一直等待，直到操作完成。不过不同于数据传输，kernel 的 launch 是异步的，host 立刻就能重新得到控制权，不用管 kernel 是否执行完毕，从而进行下一步动作。很明显，这种异步行为有助于重叠device和host之间的运算时间

数据传输也是可以异步执行的，使用时必须显示的声明一个 stream 来分派执行

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

注意：使用该函数的一个比较常见的错误，或者说容易引起混乱的地方是，这个函数返回的 error code 可能是上一次调用异步函数产生的

当执行一次异步数据传输时，必须使用pinned（或者non-pageable）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资源

1.2 stream调度

概念上来说，所有 stream 是同时运行的。但事实上并非如此

1.2.1 伪依赖

尽管 Fermi 最高支持16路并行，但是在物理上，所有 stream 是被塞进硬件上唯一一个工作队列来调度的，当选中一个 grid 来执行时，runtime 会查看 task 的依赖关系，若当前 task 依赖前面的 task，该 task 就会阻塞，由于只有一个队列，后面的都会跟着等待，即使后面的 task 是别的 stream 上的任务

C和P以及R和X是可以并行的，因为其在不同的stream中，但是ABC，PQR以及XYZ却不行。比如，在B没完成之前，C和P都在等待。

1.2.2 Hyper-Q

伪依赖的情况在 Kepler 系列里得到了解决，采用 Hyper-Q 技术。简单粗暴的理解就是，既然工作队列不够用，那就增加，于是Kepler上出现了32个工作队列。该技术也实现了 TPC 上可以同时运行 compute 和 graphic 的应用。当然，如果超过32个stream被创建了，依然会出现伪依赖的情况

1.2.3 stream优先级

对于 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);

数值较小则拥有较高优先级。若 device 不支持优先级设置，则这两个值都返回 0

1.3 使用示例

使用流实现数据传输和代码的重叠执行

cudaMemcpyAsync可以通过非默认流异步传输内存，此操作可实现内存拷贝与计算的重叠

cudaMemcpyAsync 默认情况下仅相对于主机是异步的。默认情况下，其在默认流中执行，因此对于GPU上发生的其他CUDA操作而言，是阻塞操作。但是 cudaMemcpyAsync 函数将非默认流作为可选的第5个参数。通过向其传递非默认流，内存传输可与其他非默认流中其他CUDA操作并发

#include <cstdint>
#include <iostream>
#include "helpers.cuh"
#include "encryption.cuh"void encrypt_cpu(uint64_t * data, uint64_t num_entries, uint64_t num_iters, bool parallel=true) {#pragma omp parallel for if (parallel)for (uint64_t entry = 0; entry < num_entries; entry++)data[entry] = permute64(entry, num_iters);
}__global__ 
void decrypt_gpu(uint64_t * data, uint64_t num_entries, uint64_t num_iters) {const uint64_t thrdID = blockIdx.x*blockDim.x+threadIdx.x;const uint64_t stride = blockDim.x*gridDim.x;for (uint64_t entry = thrdID; entry < num_entries; entry += stride)data[entry] = unpermute64(data[entry], num_iters);
}bool check_result_cpu(uint64_t * data, uint64_t num_entries, bool parallel=true) {uint64_t counter = 0;#pragma omp parallel for reduction(+: counter) if (parallel)for (uint64_t entry = 0; entry < num_entries; entry++)counter += data[entry] == entry;return counter == num_entries;
}int main (int argc, char * argv[]) 
{Timer timer;Timer overall;const uint64_t num_entries = 1UL << 26; // 元素数量const uint64_t num_iters = 1UL << 10; // 每个数据条目将被处理的迭代次数const bool openmp = true; // 是否并行const uint64_t num_streams = 32; // 定义流的数量// 使用四舍五入除法计算块的大小const uint64_t chunk_size = sdiv(num_entries, num_streams);timer.start();uint64_t * data_cpu, * data_gpu;cudaMallocHost(&data_cpu, sizeof(uint64_t)*num_entries);cudaMalloc    (&data_gpu, sizeof(uint64_t)*num_entries);timer.stop("allocate memory");check_last_error();timer.start();encrypt_cpu(data_cpu, num_entries, num_iters, openmp);timer.stop("encrypt data on CPU");timer.start();cudaStream_t streams[num_streams];for (uint64_t stream = 0; stream < num_streams; stream++)cudaStreamCreate(&streams[stream]);timer.stop("create streams");check_last_error();overall.start();timer.start();for (uint64_t stream = 0; stream < num_streams; stream++) {// lower索引、upper索引、数据块宽度const uint64_t lower = chunk_size * stream;const uint64_t upper = min(lower + chunk_size, num_entries);const uint64_t width = upper - lower;// 将数据块拷贝到设备cudaMemcpyAsync(data_gpu + lower, data_cpu + lower, sizeof(uint64_t) * width, cudaMemcpyHostToDevice, streams[stream]);// 流中进行计算decrypt_gpu<<<80*32, 64, 0, streams[stream]>>>(data_gpu + lower, width, num_iters);// 将数据块拷贝到主机cudaMemcpyAsync(data_cpu + lower, data_gpu + lower, sizeof(uint64_t) * width, cudaMemcpyDeviceToHost, streams[stream]);}// 同步流for (uint64_t stream = 0; stream < num_streams; stream++)cudaStreamSynchronize(streams[stream]);    timer.stop("asynchronous H2D->kernel->D2H");overall.stop("total time on GPU");check_last_error();timer.start();const bool success = check_result_cpu(data_cpu, num_entries, openmp);std::cout << "STATUS: test " << ( success ? "passed" : "failed") << std::endl;timer.stop("checking result on CPU");timer.start();for (uint64_t stream = 0; stream < num_streams; stream++)cudaStreamDestroy(streams[stream]);    timer.stop("destroy streams");check_last_error();timer.start();cudaFreeHost(data_cpu);cudaFree    (data_gpu);timer.stop("free memory");check_last_error();return 0;
}

N 能被流数量整除的情况好处理，但若是不能被整除呢？

为解决该问题，可以使用向上取整的除法运算来计算数据块大小。这样确实可以访问所有数据，但又产生了新问题：对于最后一个数据块而言，数据块大小过大

解决方法如下：

每个数据块计算 lower 索引

lower = chunk_size * streamIndex;

为每个数据块计算 upper 索引（不得超过 N ）

upper = min(lower + chunk_size, N)

使用 upper 和 lower 计算数据块 width

width = upper - lower

1.4 流同步

由于所有 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 中的操作

1.4.1 阻塞与非阻塞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 被定义成第二种，就不会被阻塞

1.4.2 隐式同步

CUDA 有两种类型的 host 和 device 之间同步：显式和隐式。如 cudaMemcpy 就会隐式的同步 device 和 host，因为该函数同步作用只是数据传输的副作用，所以称为隐式。了解这些隐式同步是很重要的，因为不经意的调用这样一个函数可能会导致性能急剧降低

若主机线程在来自不同流的两个命令之间发出以下任何一个操作，则无法同时运行这两个命令：

分页锁定的主机内存分配
设备内存分配
设备内存集
两个地址之间的内存复制到同一设备内存
任何 CUDA 命令添加到 NULL 流中
L1 / 共享内存配置之间的切换

1.4.3 显示同步

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

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

可以使用 cudaDeviceSynchronize 来同步该device上的所有操作。该函数会导致 host 等待所有 device 上的运算或者数据传输完成。该函数是个重量级的函数，应该减少这类函数的使用

通过使用 cudaStreamSynchronize 可以使 host 等待特定 stream 中的操作全部完成。或者使用非阻塞的 cudaStreamQuery 来测试是否完成

CUDA Event 可以用来实现更细粒度的阻塞与同步，相关函数为 cudaEventSynchronize 和 cudaEventQuery，用法类似 stream 相关的函数

此外，cudaStreamWaitEvent 提供了一种灵活的方式来引入 stream 之间的依赖关系

cudaError_t cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event);

该函数会指定该 stream 等待特定的 event，该 event 可以关联到相同或者不同的stream，对于不同 stream 的情况，如下图所示：

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

二、Event

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

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

Cuda API 提供了相关函数来插入 event 到 stream 中和查询该 event 是否完成。只有当该 event 标记的 stream 位置的所有操作都被执行完毕，该 event 才算完成。关联到默认 stream 上的 event 则对所有的 stream 有效

2.1 创造与销毁

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

同理 stream 的释放，在调用该函数时，若相关操作没完成，则会在操作完成后自动释放资源

2.2 记录事件和测量经过的时间

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 和 start 都关联到默认 stream 就可以了

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

2.3 可配置Events

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 只能用来同步，并且不需要记录计时数据。这样扔掉记录时间戳的消耗可以提高 cudaStreamWaitEvent 和 cudaEventQuery 的调用性能

cudaEventInterprocess 指定 event 可以被用来作为 inter-process event