一般来说,在CUDA C编程中有两个级别的并发:
(1)内核级并发
单一的内核被GPU的多个线程并行执行。
(2)网格级并发
多个内核在同一设备上同时执行。
一、流和事件概述
CUDA流是一系列异步的CUDA操作,这些操作按照主机代码确定的顺序在设备上执行。
流能封装这些操作,保持操作的顺序,允许操作在流中排队,并使它们在先前的所有操作之后执行,并且可以查询排队操作的状态。流中操作的执行相对于主机总是异步的。
在同一个CUDA流中的操作有严格的执行顺序,而在不同CUDA流中的操作在执行顺序上不受限制。使用多个流同时启动多个内核,可以实现网格级的并发。
1. CUDA流
CUDA操作(内核和数据传输)都在一个流中显示或隐式地运行。流分为:
(1)隐式声明的流(空流)
(2)显式声明的流(非空流)
如果没有显式地指定一个流,内核启动和数据传输将默认使用空流。
cudaMemcpy函数的异步版本:
cudaError_t cudaMemcpyAsync(void *dst, const void *src, size_t count, cudaMemcpyKind kind, cudaStream_t stream = 0);
在非默认流中启动内核,必须在内核执行配置中提供一个流标识符作为第四个参数:
kernel_name<<<grid, block, sharedMemSize, stream>>>(argument list);
一个非默认流声明如下:
cudaStream_t stream;
非默认流可以使用如下方式进行创建:
cudaError_t cudaStreamCreate(cudaStream_t *pStream);
cudaStreamCreate(&stream);
可以使用如下代码释放流中的资源:
cudaError_t cudaStreamDestroy(cudaStream_t stream);
在一个流中,当cudaStreamDestroy函数被调用时,如果该流中仍有未完成的工作,cudaStreamDestroy函数将立即返回,当流中所有的工作都已完成时,与流相关的资源将被自动释放。
2. 流调度
3. 流的优先级
4. CUDA事件
5. 流同步