如何理解CUDA Stream？

如何理解CUDA Stream？

参考

CUDA随笔之Stream的使用 - 知乎 (zhihu.com)

CUDA ---- Stream and Event - 苹果妖 - 博客园 (cnblogs.com)

背景

CUDA程序并发性可分为两种：

1、kernel level concurrency：一个task由GPU上多个thread并行执行的。（我们前几张介绍的都是kernel level的开发）

2、grid level level concurrency：多个task在GPU上同时执行。

（一切为了速度，冲冲冲！）

多个task并行执行会遇到一个问题，假设后续某个任务需要前面某个任务的结果，如果它们都同一时间执行，程序会崩溃的。那么，自然而然我们引出一个概念：stream。

==CUDA stream是将主机上的多个kernel进行某种顺序的排序。==

通过stream可以控制多个kernel在设备上的顺序。不同的stream互不影响。通常情况下，执行kernel比GPU-CPU传输数据的时间要长。因此，某些情况下可以将kernel的执行操作和数据传输放到不同的stream中，用kernel的时间掩盖传输时间，缩短程序运行时间。

通常情况下，执行kernel比GPU-CPU传输数据的时间要长。因此，某些情况下可以将kernel的执行操作和数据传输放到不同的stream中，用kernel的时间掩盖传输时间，缩短程序运行时间。

可以把传输数据的时间，湮灭掉。

CUDA stream可分为两种操作：

1、同步：此状态会阻塞CPU进程，直到kernel操作完成。

2、异步：此状态在唤醒kernel函数后立刻将控制权交给CPU。

还记得我们上一节介绍的SM等硬件的知识吗？这里也一样，从软件层面上看，不同的stream是同时进行的，但在硬件层面上还需要去争夺SM资源，可能需要相互等待。

GPU编程与优化实践 - 知乎 (zhihu.com)

CUDA stream可以显式或隐式调用。在前几张实例中，虽然我们在代码中没有任何stream的操作，但实际上系统会自动分配一个隐式stream，所有的kernel都在这一个stream上，如以下操作，cudaMemcpy会阻塞cpu进程，直到数据传输操作完成。

但假如我们相对多个kernel的执行顺序进行操作，则必须申请一个或多个显式stream进行管理。例如在以下情形中：

1、重叠host和device的计算。

2、重叠host计算与cpu-gpu的数据传输。

3、重叠device计算与cpu-gpu的数据传输。

4、gpu的并发计算

cudaMemcpyAsync可以执行异步操作，在数据传输过程中，将操作权交给cpu进行控制。

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

申请cuda stream：

cudaError_t cudaStreamCreate(cudaStream_t* pStream);

在异步传输数据时必须使用固定/不可分页的主机内存，可使用以下APT申请：

cudaError_t cudaMallocHost(void **ptr, size_t size);
cudaError_t cudaHostAlloc(void **pHost, size_t size, unsigned int flags);

CUDA Events是cuda stream中的标记。（类似跑马拉松，把整个路程分为一段段的，每一段的路标就是一个events）Event只是一个标志，你可以用来查询stream的状态：

1、同步stream的执行

2、监控device进度

CUDA允许在任何一个点插入event，并查询该event是否完成。只有当该事件段中所有操作都被满足，event才会标志为完成。

events的声明，创建，销毁：

cudaEvent_t event;

cudaError_t cudaEventCreate(cudaEvent_t* event);

cudaError_t cudaEventDestroy(cudaEvent_t event);

---

CUDA ---- Stream and Event - 苹果妖 - 博客园 (cnblogs.com)

这个博客写的也非常系统：

---

请在这里，系统地学习CUDA编程

CUDA Zone - Library of Resources | NVIDIA Developer

CUDA Toolkit Documentation 12.2 (nvidia.com)

---

CUDA优化方案—stream的使用_cudastream_t_A晨的博客的博客-CSDN博客

CUDA Stream的高阶用法：

//创建多个流
cudaStream_t stream[3];
for(int i=0; i<3; i++)
    cudaStreamCreat(&stream[i]);

float* hostPtr;
cudaMallocHost(&hostPtr, 3*size);
......
//开启了三个流，将数据分成三份分别交给了三个流进行计算
for(int i=0; i<3; i++){
    //从CPU内存复制数据到GPU显存
    cudaMemcpyAsync(DevPtr + i*size, hostPtr + i*size, size, cudaMemcpyHostToDevice, stream[i]);
    //执行Kernel函数
    Mykernel<<<grid, block, 0, stream[i]>>>(outputDevPtr + i*size, size, cudaMemcpyDeviceToHost, stream[i]);
    //将结果拷贝回CPU内存
    cudaMemcpyAsync(hostPtr + i*size, outputDevPtr + i*size, size, cudaMemcpyDeviceToHost, stream[i]);
}
//同步流
for(int i=0; i<3; i++)
    cudaStreamSynchronise(stream[i]);
//销毁流
for(int i=0; k<3; i++)
    cudaStreamDestroy(stream[i]);

}