并行编程实战——CUDA编程的Stream

一、CUDA中的Stream

提到流（Stream），大家可能非常熟悉，如学习计算机技术时遇到的指令流、命令流和数据流等等流。同样在CUDA中也存在流这个概念，那么在CUDA中流是什么意思呢？流就是CUDA中的一个操作队列（和CPU中的流类似）。
CUDA中流就是一个由主机控制的次序的执行队列，如果有多个流，则如CPU中的并行指令流一样，同样无法确定他们间的执行顺序（可能顺序也可能交错亦可能并发）。CUDA中的所有操作都必须通过流来实现，流可以分成两种类型即默认流（空流）和非默认流。
从另外一个角度，流又可以分为阻塞流和非阻塞流。相对主机来说，所有的非默认流的操作都是非阻塞的（异步），而默认流则为阻塞的（同步，但在CUDA 7后，可以通过编译选项 -default-stream per-thread 编译选项或定义 CUDA_API_PER_THREAD_DEFAULT_STREAM宏来处理流的操作方式）。同步流一般会阻塞主机的操作但不包括kernel launch的操作。常见cudaStreamCreate函数，创建的就是阻塞流。

二、流的作用和多流

在前面的分析中，大家可以知道，GPU最强大之处在于并行，而GPU并行一般来说分成两部分即：Kernel层级和Grid层级。大家熟悉的Thread和task一般指的是Kernel层级的。而Stream则指的是Grid层级的，通过创建多个流（非默认流），实现并行执行。即流的最重要的作用就是实现核函数的外部并行，它能够有效的加速程序的并行程序从而提高效率。

三、流的调度

在流的介绍中提到了，流是在Grid层级的并行。可能很多人会以为其天然是并行的，但实际情况并非这样。比如硬件提供的执行队列是有限的，而任务流的创建可以是无限的。如果创建的任务互相依赖导致所有的队列无法执行下这下，这就会出现CUDA中的伪依赖（False Dependencies）。
为了解决伪依赖，最简单的方法就是增加执行队列（Hyper-Q），但从物理的角度看，其相对软件的任务总是不足。所以，在较高版本的CUDA中，在Stream中增加了Priorities。通过优先级的处理，尽可能打破了任务的依赖，这个和CPU的任务处理类似，应该非常容易理解。
是不是想到了流水线，对，在GPU中也是有流水线的概念的，使用流可以进行相关的流水线的处理。

四、流和事件

在流的应用中，可以通过cudaStreamSynchronize和cudaDeviceSynchronize函数实现流的简单同步。但实际的应用往往需要更精细化的同步处理，这时事件就登场了。
事件，作为同步机制，有经验的开发者应该都听说过。CUDA中的事件与其类似，它是在流中创建标记点，刻录相关的的操作时间，从而将流操作可确定化。CUDA接口中提供了cudaEventCreate函数用来创建事件。事件的主要功能有：测量时间，同步流和管理资源（通过事件来处理各种CUDA操作已经完成后再回收相关资源）。

五、流和图

正如并行的执行算法往往会应用到图的情况，CUDA中也有相关的CUDA图。一个操作可以形成一个节点，操作间的依赖关系可以认为是边，通过这些节点和边的联系就可以确定操作间的依赖执行顺序即可执行图，从而由CUDA确定进行调度。
可执行图可以加载到流中，这与其它的操作行为在CUDA中的工作无异。它可以在不重复实例化的情况下反复启动进行操作。同样，可以通过流捕获创建图即从现有的基于流的API创建图的机制。其实也可以叫做图和流的绑定。开发者需要注意的是，要处理好流退出捕获模式下的依赖关系的处理以及异常错误的处理。细节可查看CUDA的相关文档。

六、CUDA中流的应用

CUDA中流的作用非常重要，下面看一下相关的应用：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

__device__ __managed__ int x, y = 2;
__global__ void kernel(int* ptr) {
	//cuda::atomic_ref{ *ptr }.store(2);
}
int main()
{

	cudaStream_t stream1, stream2;
	cudaStreamCreate(&stream1);
	cudaStreamCreate(&stream2);
	int *non_managed, *managed, *also_managed;
	cudaMallocHost(&non_managed, 4);    // Non-managed, CPU-accessible memory
	cudaMallocManaged(&managed, 4);
	cudaMallocManaged(&also_managed, 4);
	// Point 1: CPU can access non-managed data.
	kernel << < 1, 1, 0, stream1 >> >(managed);
	*non_managed = 1;
	// Point 2: CPU cannot access any managed data while GPU is busy,
	//          unless concurrentManagedAccess = 1
	// Note we have not yet synchronized, so "kernel" is still active.
	*also_managed = 2;      // Will issue segmentation fault
							// Point 3: Concurrent GPU kernels can access the same data.
	kernel << < 1, 1, 0, stream2 >> >(managed);
	// Point 4: Multi-GPU concurrent access is also permitted.
	cudaSetDevice(1);
	kernel << < 1, 1 >> >(managed);

    // cudaDeviceReset must be called before exiting in order for profiling and
    // tracing tools such as Nsight and Visual Profiler to show complete traces.
    cudaDeviceReset();

    return 0;
}