流和并发

流和并发

一、流和事件概述

CUDA流是一系列异步的CUDA操作，这些操作按照主机代码确定的顺序在设备上执行。流能够封装这些操作，保持操作的顺序，允许操作在流中排队，并使它们在先前的操作之后执行，并且可以查询排队操作的状态。这些操作包括在主机和设备间进行数据传输，内核启动以及大多数由主机发起但由设备执行的其他命令。流中操作的执行相对于主机总是异步的。CUDA运行时决定何时可以在设备上执行操作。使用流可以充分利用GPU和CPU的闲置资源（包括计算资源和内存带宽）。

一般来说，CUDA典型的编程模式为：

1.将数据从主机复制到设备

2.执行核函数

3.将数据从设备复制到主机

一般来看，这个过程必须按顺序执行，但事实上，执行1时，计算资源和到主机的pcie带宽浪费，执行2时，两个pcie带宽浪费，执行3时计算资源和主机到设备的pcie带宽浪费。而这些资源浪费时可以通过将内核调度到不同的流中，使操作重叠来减少浪费，更加充分的利用GPU资源。

从软件看来，不同的流中并发是完全并行的，但从硬件来看，不一定总是并行的，根据PCIe总线占用或SM资源的可用性，完成不同的CUDA流可能仍需要互相等待。

1.1 CUDA流

流的两种类型：

隐式声明的流（空流）

显示声明的流（非空流）

在之前的操作里，并没有提到流的概念，是因为，所有的操作都被加载到空流中了。

非空流可以被显式的创建和管理。想要重叠不同的CUDA操作，需要使用非空流。异步流可以支持以下几种粗粒度并发：

1.重叠主机计算和设备计算

2.重叠主机计算和主机与设备间的数据传输

3.重叠主机与设备间的数据传输和设备计算

4.并发设备计算

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

是cudaMemcpy的异步形式（调用之后立刻返回），最后一个参数为流标识符，默认情况下为空流。

cudaStream_t stream;//定义一个流标识符

cudaError_t cudaStreamCreate(cudaStream_t *pStream);//创建一个可以显式管理的流

。。。使用一个流。一般api函数中流有其自己的形参，核函数的话，将流放在核函数执行配置的第四个参数即可（第三个是动态申请的共享内存的大小）

cudaError_t cudaStreamDestroy(cudaStream_t stream);//销毁一个流当执行资源释放的时候，如果仍然有stream的工作没干完，那么虽然该函数仍然会立刻返回，但是相关的工作做完后，这些资源才会自动的释放掉。

使用异步数据传输时，必须使用固定主机内存（因为主机端已经无法知道什么时候这些内存会被拷贝，分页内存会导致未定义行为）。即使用cudaMallocHost函数或cudaHostAlloc函数分配固定内存。

使用异步函数时，该函数没有执行完就已经返回，因此可能导致无法正确探查执行过程中的错误，或返回的错误是先前的异步操作的错误代码。

cudaError_t cudaStreamSynchronize(cudaStream_t stream);

cudaError_t cudaStreamQuery(cudaStream_t stream);

cudaStreamSynchronize强制阻塞主机，知道给定流中所有操作都完成

cudaStreamQuery会检查流中所有操作是否都完成，但不会阻塞主机。当所有操作都完成时，返回cudaSuccess，有任务还在执行时返回cudaErrorNotReady。

一个比较简单的例子就是cudaSamples里面的simpleStream，广度优先方式（先将一个任务分的段分别加载到各个流中，然后将下一个任务的各个段再加载到任务中）。

异步方式：

// asynchronously launch nstreams kernels, each operating on its own portion of data
for (int i = 0; i < nstreams; i++)
{
  init_array<<<blocks, threads, 0, streams[i]>>>(d_a + i *n / nstreams, d_c, niterations);
}
// asynchronously launch nstreams memcopies.  Note that memcopy in stream x will only
//   commence executing when all previous CUDA calls in stream x have completed
for (int i = 0; i < nstreams; i++)
{
  checkCudaErrors(cudaMemcpyAsync(hAligned_a + i * n / nstreams, d_a + i * n / nstreams, nbytes / nstreams, cudaMemcpyDeviceToHost, streams[i]));
}

结果如下图所示。

同步方式：

init_array<<<blocks, threads>>>(d_a, d_c, niterations);//这里的blocks和threads和上面代码不一样，维度分别应该是上面的nstreams倍
checkCudaErrors(cudaMemcpy(hAligned_a, d_a, nbytes, cudaMemcpyDeviceToHost));

可以看出核函数和D2H方向的数据拷贝在时间轴上是重叠的，达到了我们的目的。比下图中默认流的方式要好。

某些PCIe总线是双工的，可以重叠两个不同流中且不同方向的数据传输

并发内核的最大数量依赖于设备。Fermi支持16路并发，Kepler支持32路并发。而当前可用的计算资源如计算单元、寄存器、共享内存等都会限制并发内核的数量，如果当前设备是满负荷的，那么不能进一步并发。

1.2 流调度

Hyper-Q技术。

从Kepler架构开始有。目的是为了减少虚假的依赖关系。Hyper-Q技术通过在主机和设备之间维持多个硬件管理上的链接，允许多个CPU线程或进程在单一GPU上同时启动工作（其实没懂什么意思，但现在我们用的Maxwell，Pascal应该都支持这种技术）。

流的优先级。

从计算能力3.5开始有。使用

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

创建一个有指定整数优先级的流。并在pStream中返回一个句柄，用于工作调度。优先级整数越小，越优先。流优先级不会影响数据传输操作，只对计算内核有影响。高优先级网格队列可以占用低优先级已经执行的工作。如果给定整数超过设备定义范围，则被限制为最大值或最小值。可用下函数查询优先级

cudaError_t cudaDeviceGetStreamPriorityRange(int leastPriority, int greatestPriority);

若不支持流优先级返回两个0.

CUDA事件。

本质是CUDA流中的标记，它与该流内操作流中特定点相关联。不涉及。

流同步（阻塞和非阻塞）。

使用cudaStreamCreate创建的流是阻塞流，这些流的操作被会被空流阻塞。即考虑如下情况：

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

kernel_1和kernel_3被kernel_2阻塞，所以kernel_1执行完才会执行kernel_2，kernel_2执行完才能执行kernel_3，但对于主机，仍然是所有都是异步的，对于设备则是被阻塞的。

使用下面函数可以让流成为非阻塞流

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

flags参数使用cudaStreamNonBlocking即可。

隐式同步。

显式调用同步函数为显式同步。隐式同步为调用类似cudaMemcpy的函数，造成同步。

可能造成隐式同步的操作有：

锁页主机内存分配

设备内存分配

设备内存初始化

同一设备上两个地址之间的内存复制

一级缓存/共享内存配置的修改

显式同步。

cudaDeviceSynchronize、cudaStreamSynchronize、cudaStreamQuery、cudaEventSynchronize、cudaEventquery、cudaStreamWaitEvent等函数。

可配置事件。

不涉及。

2.并发内核执行

Fermi架构中存在虚假依赖关系（一直没懂是什么，反正见不到这种设备了），需要使用广度优先方法避免（深度优先比广度优先慢很多）。而又Hyper-Q技术的架构则深度优先和广度优先速度差不多。不多探讨了。

cudaSamples中simpleHyperQ例子，结果如下（有八个流能够同时执行）

使用OpenMP的调度操作

略

用环境变量调整流行为

使用CUDA_DEVICE_MAX_CONNECTIONS环境变量来调整并行硬件连接数量（对于Kepler架构最大为32），由于每个kernel分配的资源多少不同，实际可并发的可能更少。

建立流间的依赖关系

由于实际情况复杂，引入流的依赖关系可能是有用的。

for (int i = 0; i < n_streams; i++)
{
    CHECK(cudaEventCreateWithFlags(&(kernelEvent[i]),
                cudaEventDisableTiming));
}
​
// record start event
CHECK(cudaEventRecord(start, 0));
​
// dispatch job with depth first ordering
for (int i = 0; i < n_streams; i++)
{
    kernel_1<<<grid, block, 0, streams[i]>>>();
    kernel_2<<<grid, block, 0, streams[i]>>>();
    kernel_3<<<grid, block, 0, streams[i]>>>();
    kernel_4<<<grid, block, 0, streams[i]>>>();
​
    CHECK(cudaEventRecord(kernelEvent[i], streams[i]));
    CHECK(cudaStreamWaitEvent(streams[n_streams - 1], kernelEvent[i], 0));
}

注意第n_streams-1个流需要等待其他流完成才能开始启动工作。

3 重叠内核执行和数据传输

实际上就是1.1中的例子。同步方式和异步方式的对比。1.1就是将1个完整的复制、计算、复制，分块成n个小块的复制1、计算1、复制1，复制2、计算2、复制2。。。复制n、计算n、复制n。实现重叠。提高资源利用。

将一个大的问题拆分成相同的子问题，然后通信和计算就可以重叠了。

4 重叠CPU和GPU的执行

异步方式执行设备操作，使用cudaEventRecord记录一个停止事件，然后执行主机端操作，需要设备数据时使用cudaEventquery查询停止事件是否执行完毕。详细例子如cudaSamples中asyncAPI。

5 流回调

看起来挺高级的功能，暂时不用，略。

注意：

使用nvvp可视化查看cuda操作的时间轴重叠对于本节实验很重要。

流只在GPU有资源可用的时候才能起作用。比如计算满负荷时，内存拷贝是可以执行的，但新的计算时不能被执行的。双工PCIe的D2H和H2D是可以被重叠的。
参考：CUDA C编程权威指南，第六章cudaSamples asyncAPI、simpleHyperQ、simplestreams