CUDA进阶：流与事件

流

CUDA 流是一系列异步的 CUDA 操作,这些操作按照主机代码确定的顺序在设备上执行。流能封装这些操作,保持操作的顺序,允许操作在流中排队,并使它们在先前的所有操作之后执行,并且可以查询排队操作的状态。这些操作包括在主机与设备间进行数据传输,内核启动以及大多数由主机发起但由设备处理的其他命令。流中操作的执行相对于主机总是异步的。 CUDA 运行时决定何时可以在设备上执行操作。我们的任务是使用 CUDA 的 API 来确保一个异步操作在运行结果被使用之前可以完成。

在同一个 CUDA 流中的操作有严格的执行顺序,而在不同CUDA 流中的操作在执行顺序上不受限制。使用多个流同时启动多个内核,可以实现网格级并发。

CUDA 的 API 函数一般可以分为同步或异步。具有同步行为的函数会阻塞主机端线程,直到它们完成。具有异步行为的函数被调用后,会立即将控制权归还给主机。异步函数和流是在 CUDA 中构建网格级并发的两个基本支柱。

• 基于流的异步内核启动和数据传输支持以下类型的粗粒度并发

  1. 重叠主机和设备计算
  2. 重叠主机计算和主机设备数据传输
  3. 重叠主机设备数据传输和设备计算
  4. 并发设备计算(包括多个设备)

• 不支持并发

  • a page-locked host memory allocation,
    主机锁页内存分配，
  • a device memory allocation,
    设备内存分配，
  • a device memory set,
    设备内存设置
  • a memory copy between two addresses to the same device memory,
    相同设备两块内存区域间的内存拷贝
  • any CUDA command to the NULL stream
    将CUDA命令加入到默认流
  • L1/共享内存配置的修改

NULL流会强制所有的GPU引擎汇聚，所以如果指定了NULL流，那么即使一个异步的内存复制也会使多个引擎之间的并发停止。

流主要应用场景：

（1）计算和传输同时重叠进行（主要适用于独立显卡）。

（2）多流中的小kernel累加起来，充分利用设备。

异步的内存复制

像内核启动一样，异步的内存复制调用在GPU完成待处理的内存复制之前就会返回。由于GPU能够自主运行，并且可以在没有任何操作系统接入。只有锁页内存有资格进行异步的内存复制。

异步内存复制函数都用Async()作为后缀。

锁页内存申请 cudaHostAlloc( (void**) p , size ,cudaHostAllocDefault)

//声明
cudaStream_t stream0,stream1;
// 创建一个异步流
//__host__ cudaError_t cudaStreamCreate ( cudaStream_t* pStream )
cudaStreamCreate( &stream0 );
cudaStreamCreate(