CUDA通过流来实现网格级并发。
流和事件
CUDA流是一系列异步的CUDA操作，这些操作按照主机代码确定的顺序在设备上执行。流可以封装这些操作，保持操作的顺序，允许操作在流中排队，并使他们在先前的所有操作之后执行。
这些操作包括在主机与设备之间进行数据传输，内核启动以及大多数由主机发起但由设备处理的其他命令。
流中操作的执行相对于主机总是异步的。CUDA运行时决定何时可以在设备上执行操作。我们的任务是使用CUDA的API来确保一个异步操作在运行结果被使用之前可以完成。
在同一个CUDA流中的操作有严格的执行顺序，而在不同的CUDA流中的操作在执行顺序上不受限制。使用多个流同时启动多个内核，可以实现网格级并发。
CUDA编程的一个典型模式是：
- 数据从主机移到设备上
- 设备上执行一个内核
- 将结果从设备移回主机
在多数情况下，执行内核比传输数据耗时的多。在这些情况下，可以完全隐藏CPU和GPU之间的通信延迟。将内核执行和数据传输调度到不同的流中，这些操作可以重叠，程序的总运行时间将被缩短。流在CUDA的API调用粒度上可实现流水线或双缓冲技术。
CUDA的api函数分为同步和异步。同步函数会阻塞主机端线程，直到其完成。异步函数被调用后，控制权直接移回主机。异步函数和流是在CUDA中构建网格级并发的两个基本支柱。
CUDA流
所有的CUDA操作（内核和数据传输）都在一个流中显式或隐式的运行。流分为两种：
- 隐式声明的流（空流）
- 显式声明的流（非空流）
如果没有显式的声明一个流，那么内核启动和数据传输将默认使用空流。
非空流可以被显式的创建和管理。如果想要重叠不同的CUDA操作，必须使用非空流。基于流的异步的内核启动和数据传输支持以下类型的并发：
- 重叠主机计算和设备计算
- 重叠主机计算和主机与设备间的数据传输
- 重叠主机与设备间的数据传输和设备计算
- 并发设备计算
思考下面使用默认流的方法：
cudaMemcpy(...,cudaMemcpyHostToDevice);
kernel<<<grid,block>>>();
cudaMemcpy(...,cudaMemcpyDeviceToHost);从设备角度来看，上述代码的所有3个操作都被发布到默认流中，并且按发布顺序执行。设备不知道其他被执行的主机操作。
从主机角度看，数据传输是同步的，强制空闲主机等待数据传输完成。内核启动是异步的，无论内核是否完成，主机的应用程序都立即恢复执行。这种内核启动的默认异步行为使它可以直接重叠设备和主机计算。
数据传输也可以异步发布，但是必须显式的设置一个CUDA流来装载。提供以下函数
cudaError_t cudaMemcpyAsync(void* dst, void* src, size_t count, cudaMemcpyKind kind, cudaStream_t stream=0);附加的流的标识符作为第五个参数，默认情况下使用默认流。这个函数与主机是异步的，所以调用发布后，控制权将立即返回到主机。
使用如下代码创建一个非空流
cudaError_t cudaStreamCreate(cudaStream_t* pStream);该函数创建了一个可以显示管理的非空流。之后，返回到pStream中的流就可以被当做流参数供其他的异步CUDA的API函数使用。
在执行异步数据传输时，必须使用固定主机内存。在非默认流中启动内核，必须在内核执行配置中提供一个流标识符作为第四个参数
kernel<<<grid, block, shareMemSize, stream>>>()非默认流的声明和创建如下：
cudaStream_t stream;
cudaStreamCreate(&stream);
//释放资源
cudaStreamDestroy(cudaStream_t stream);在一个流中，当cudaStreamDestroy被调用时，如果该流中仍有未完成的工作，函数将立即返回，当流中所有工作都已完成时，与流相关的资源将被自动释放。
用以下两个函数检查流中工作是否完成
cudaError_t cudaStreamSynchronize(cudaStream_t stream);
cudaError_t cudaStreamQuery(cudaStream_t stream);synchronize将阻塞主机，直到给定流中的工作全部完成。query检查流中所有操作是否完成，但不会阻塞，如果完成返回cudaSuccess，未完成返回cudaErrorNotReady。
画个图展示流的作用
for(int i=0; i<nStreams; i++){
int offset = i * bytesPerStream;
cudaMemcpyAsync(&d_a[offset], &a[offset], bytesPerStream, 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]);
}
并发内核的最大数量是依赖设备的。费米架构支持16路并发，开普勒设备支持32路并发。
流调度
从概念上讲，所有流可以并发执行。但是，当流映射到物理硬件时，并不总是这样的。
虚假的依赖关系
虚假的依赖关系（False Dependency）指的是一种由编译器或硬件引入的假象性依赖，导致了代码中不必要的序列化或延迟。这种依赖关系并不反映真实的数据依赖关系，但会影响到代码的执行顺序和性能。
虚假的依赖关系通常出现在对共享内存的操作中，特别是在使用指针进行多次访存的情况下。编译器或硬件可能会认为对同一内存地址的多次访问之间存在依赖关系，从而引入不必要的序列化或延迟。
Hyper-Q技术
Hyper-Q技术使用多个硬件工作队列，减少了虚假的依赖关系。允许多个CPU线程或进程在单一GPU上同时启动工作。
开普勒GPU使用32个工作队列，每个流分配一个工作队列。如果超过32个流，多个流将共享一个硬件工作队列。
流的优先级
计算能力3.5或更高的设备，可以给流分配优先级。优先级高的流的网格队列可以优先占有低优先级流的已经执行的工作。
CUDA事件
CUDA中事件本质上是CUDA流中的标记，它与流内操作中特定点相关联。可以使用事件来执行以下两个基本任务：
- 同步流的执行
- 监控设备的进展
CUDA的API提供了在流中任意点插入事件以及查询事件完成的函数。只有当一个给定CUDA流中先前的所有操作都执行结束后，记录在该流内的事件才会起作用（即完成）。
创建和销毁
声明
cudaEvent_t event;创建
cudaError_t cudaEventCreate(cudaEvent_t* event);销毁
cudaError_t cudaEventDestroy(cudaEvent event);记录事件和计算运行时间
事件在流中标记了一个点。可以用来检查正在执行的流操作是否已到达了给定点。
使用如下函数进入CUDA流
cudaError_t cudaEventRecord(cudaEvent_t event, cudaStream_t stream=0);使用如下函数计算两个事件标记的CUDA操作运行时间
cudaError_t cudaEventElapsedTime(float* ms, cudaEvent_t start, cudaEvent_t end);返回运行时间，单位为ms。
示例如何进行计时
//创建两个事件
cudaEvent_t start, end;
cudaEventCreate(&start);
cudaEventCreate(&end);
//将开始事件记录在默认流中
cudaEventRecord(start);
kernel<<<>>>();
//将结束事件记录在默认流中
cudaEventRecord(end);
//等待事件结束
cudaEventSynchronize(end);
//计算时间
float time;
cudaEventElapsedTime(&time, start, end);
//释放事件
cudaEventDestroy(start);
cudaEventDestroy(end);cudaEventSynchronize和stream的相关函数相同。
CUDA流中的事件可以在主机端和设备端都记录。在CUDA中，事件（Event）用于测量时间间隔或同步CUDA流中的操作。主要有以下两种类型的事件：
主机事件（Host Event）：主机事件是由主机代码创建和记录的事件，用于测量主机和设备之间的时间间隔或同步主机代码和CUDA流中的操作。可以使用
cudaEventRecord()函数记录主机事件。设备事件（Device Event）：设备事件是由设备代码（即在CUDA核函数中）创建和记录的事件，用于测量CUDA流中的操作的时间间隔或同步不同的CUDA核函数。可以使用
cudaEventRecord()函数在设备代码中记录设备事件。在使用CUDA流时，通常会在主机端记录主机事件来测量主机与设备之间的时间间隔，同时也可以在设备端使用设备事件来测量CUDA核函数的执行时间或同步不同核函数之间的操作。
