CUDA从入门到放弃（七）：流（ Streams）
应用程序通过流来管理并发操作，流是一系列按顺序执行的命令。不同的流可能无序或并发地执行命令，但此行为并不保证。流上的命令在依赖关系满足时执行，这些依赖可能来自同一流或其他流。同步调用（synchronize call）可以确保所有启动的命令已完成。
任何 CUDA 操作都存在于某个 CUDA 流中，要么是默认流（default stream），也称为空流（null stream），要么是明确指定的非空流。
1 流的基本概念
1-1 流的创建与销毁 Creation and Destruction of Streams
CUDA 流的定义、产生与销毁
cudaStream_t stream_1;
cudaStreamCreate(&stream_1); // 注意要传流的地址
cudaStreamDestroy(stream_1);
示例：
cudaStream_t stream[2];
for (int i = 0; i < 2; ++i)
cudaStreamCreate(&stream[i]);
float* hostPtr;
cudaMallocHost(&hostPtr, 2 * size);
for (int i = 0; i < 2; ++i) {
cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size,
size, cudaMemcpyHostToDevice, stream[i]);
MyKernel <<<100, 512, 0, stream[i]>>>
(outputDevPtr + i * size, inputDevPtr + i * size, size);
cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size,
size, cudaMemcpyDeviceToHost, stream[i]);
}
for (int i = 0; i < 2; ++i)
cudaStreamDestroy(stream[i]);
1-2 默认流 Default Stream
my_kernel<<<N_grid, N_block>>>(函数参数);
my_kernel<<<N_grid, N_block, N_shared>>>(函数参数);
my_kernel<<<N_grid, N_block, N_shared, stream_id>>>(函数参数);
如果用第一种调用方式，说明核函数没有使用动态共享内存，而且在默认流中执行；
如果用第二种调用方式，说明核函数在默认流中执行，但使用了 N_shared 字节的动态共享内存；
如果用第三种调用方式，则说明核函数在编号为stream_id 的 CUDA 流中执行，而且使用了 N_shared 字节的动态共享内存。
在使用非空流但不使用动态共享内存的情况下，必须使用上述第三种调用方式，并将 N_shared 设置为零：
my_kernel<<<N_grid, N_block, 0, stream_id>>>(函数参数); // 正确
1-3 流状态查询
为了实现不同 CUDA 流之间的并发，主机在向某个 CUDA 流中发布一系列命令之后必须马上获得程序的控制权，不用等待该 CUDA 流中的命令在设备中执行完毕。这样，就可以通过主机产生多个相互独立的 CUDA 流。
为了检查一个 CUDA 流中的所有操作是否都在设备中执行完毕， CUDA 运行时 API 提
供了如下函数：
__host__cudaError_t cudaLaunchHostFunc(cudaStream_t stream, cudaHostFn_t fn, void *userData)
cudaDeviceSynchronize()会等待直到所有主机线程的所有流中的所有前面命令都已完成。
__host__cudaError_t cudaStreamSynchronize(cudaStream_t stream)
cudaStreamSynchronize()接受一个流作为参数，并等待直到给定流中的所有前面命令都已完成。它可用于同步主机与特定流，同时允许其他流继续在设备上执行。
__host____device__cudaError_t cudaStreamWaitEvent (cudaStream_t stream, cudaEvent_t event, unsigned int flags)
cudaStreamWaitEvent()接受一个流和一个事件作为参数，并使得所有在调用cudaStreamWaitEvent()之后添加到给定流的命令延迟其执行，直到给定事件已完成。
cudaError_t cudaStreamQuery(cudaStream_t stream);
函数 cudaStreamQuery 不会阻塞主机，只是检查 CUDA 流 stream 中的所有操作是否都执行完毕。若是，返回 cudaSuccess，否则返回 cudaErrorNotReady。
2 重叠执行 Overlapping Behavior
2-1 在默认流中重叠主机和设备计算
虽然同一个 CUDA 流中的所有 CUDA 操作都是顺序执行的，但依然可以在默认流中重叠主机和设备的计算。
示例：
cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);
cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice);
sum<<<grid_size, block_size>>>(d_x, d_y, d_z, N);
cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);
从
cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);
到
sum<<<grid_size, block_size>>>(d_x, d_y, d_z, N);
顺序依次执行。
但是当主机发出
sum<<<grid_size, block_size>>>(d_x, d_y, d_z, N);
命令之后，不会等待该命令执行完毕，而会立刻得到程序的控制权。主机紧接着会发出从设备到主机传输数据的命令
cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);
该命令会在CUDA流中等待前一个 CUDA 操作（即核函数的调用）执行完毕才会开始执行。
因此，主机在发出核函数调用的命令之后，可以执行主机中的某个计算任务，那么主机就会在设备执行核函数的同时去进行一些计算。这样，主机和设备就可以同时进行计算。
2-2 用非默认 CUDA 流重叠多个核函数的执行与数据传递
尽管在单个默认流中确实可以实现主机计算与设备计算的并行处理，但要实现多个核函数之间的并行执行，则必须借助多个CUDA流。这是因为，在相同的CUDA流内，CUDA操作在GPU设备上是按照顺序执行的。因此，同一个CUDA流内的核函数也必须在设备上依次执行，即便主机在发出每一个核函数调用指令后都立即恢复了程序的控制权。这样的设计确保了操作的顺序性和一致性，但同时也限制了并行度的提升。通过使用多个CUDA流可以有效地将不同的核函数任务分配到不同的执行流中，从而实现更高程度的并行计算，提升整体计算性能。
为了实现核函数执行与数据传输的并发（重叠），必须确保这两个操作在不同的非默认流中执行，且数据传输应使用 cudaMemcpyAsync 函数，这是 cudaMemcpy 的异步版本。cudaMemcpyAsync 通过 GPU 中的 DMA 实现直接内存访问，无需主机参与。
异步传输函数 cudaMemcpyAsync：
cudaError_t cudaMemcpyAsync(
void *dst, // 目标内存地址
const void *src, // 源内存地址
size_t count, // 传输的字节数
enum cudaMemcpyKind kind, // 传输方向（主机到设备、设备到主机等）
cudaStream_t stream // 所在的流
);
这个函数与同步版本的 cudaMemcpy 相比，仅多出一个参数：流（stream）。通过将数据传输操作分配给特定流，并允许核函数在另一个流中执行，可以实现数据传输和核函数执行的并发执行，从而提高程序的整体性能。
示例：
for (int i = 0; i < 2; ++i)
cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size,
size, cudaMemcpyHostToDevice, stream[i]);
for (int i = 0; i < 2; ++i)
MyKernel<<<100, 512, 0, stream[i]>>>
(outputDevPtr + i * size, inputDevPtr + i * size, size);
for (int i = 0; i < 2; ++i)
cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size,
size, cudaMemcpyDeviceToHost, stream[i]);
3 主机函数（回调函数）Host Functions (Callbacks)
通过 cudaLaunchHostFunc()，可以在流的任意位置插入主机函数的调用。此函数会在流中所有先前命令完成后在主机上执行。
以下示例在每个流执行主机到设备内存复制、内核启动和设备到主机内存复制后，将主机函数 MyCallback 添加到流中。一旦设备到主机的内存复制完成，该回调函数将在主机上执行。
void CUDART_CB MyCallback(void *data){
printf("Inside callback %d\n", (size_t)data);
}
...
for (size_t i = 0; i < 2; ++i) {
cudaMemcpyAsync(devPtrIn[i], hostPtr[i], size, cudaMemcpyHostToDevice, stream[i]);
MyKernel<<<100, 512, 0, stream[i]>>>(devPtrOut[i], devPtrIn[i], size);
cudaMemcpyAsync(hostPtr[i], devPtrOut[i], size, cudaMemcpyDeviceToHost, stream[i]);
cudaLaunchHostFunc(stream[i], MyCallback, (void*)i);
}
请注意，排入流的主机函数应避免直接或间接调用 CUDA API，以免发生死锁。
4 流优先级
在创建流时，可以使用 cudaStreamCreateWithPriority() 函数指定流的相对优先级。通过 cudaDeviceGetStreamPriorityRange() 函数，可以获取允许的优先级范围，该范围按从最高优先级到最低优先级的顺序排列。在运行时，高优先级流中的待处理任务将优先于低优先级流中的待处理任务执行。
以下代码示例获取当前设备的允许优先级范围，并使用可用的最高和最低优先级创建流。
// get the range of stream priorities for this device
int priority_high, priority_low;
cudaDeviceGetStreamPriorityRange(&priority_low, &priority_high);
// create streams with highest and lowest available priorities
cudaStream_t st_high, st_low;
cudaStreamCreateWithPriority(&st_high, cudaStreamNonBlocking, priority_high);
cudaStreamCreateWithPriority(&st_low, cudaStreamNonBlocking, priority_low);
参考资料
1 CUDA编程入门
2 CUDA编程入门极简教程
3 CUDA C++ Programming Guide
4 CUDA C++ Best Practices Guide
5 NVIDIA CUDA初级教程视频
6 CUDA专家手册 [GPU编程权威指南]
7 CUDA并行程序设计：GPU编程指南
8 CUDA C编程权威指南