流（stream）

设备操作包括：数据传输和执行kernel函数。

在cuda中，所有的设备操作都在stream中执行。当没有指定stream时，使用默认的stream。

默认stream

默认stream是一个针对设备操作同步的stream，也就是说，只有当所有之前设备上任何stream(包括默认stream)里面的操作全部完成时，才开始默认stream里面操作的执行，并且默认stream里面的一个操作必须完成，其他任何stream（包括默认stream）里面的操作才能开始。

举个例子：

cudaMemcpy(d_a, a, numBytes, cudaMemcpyHostToDevice); 
increment<<<1,N>>>(d_a);
cudaMemcpy(a, d_a, numBytes, cudaMemcpyDeviceToHost);
    1. 从设备端来看，这三个操作都在默认stream中，并且按顺序执行；
    1. 从主机端来看，数据传输是阻塞的或者同步传输，而kernel是异步的；
    • 第一步主机到设备的数据传输是同步的，CPU线程不能到达第二行直到主机到设备的数据传输完成。
    • 一旦kernel被launch，CPU线程移到第三行，但是该行的传输还不能开始，因为设备端正在执行第二行的内容（但如果在2-3行之间插入其他cpu操作是可以执行的）。

非默认stream

    1. 非默认stream中的数据传输使用函数cudaMemcpyAsync()，这个函数在主机端是非阻塞的，传输处理后控制权马上返回给主机线程；
    1. 自定义一个stream可以用cudaError_t cudaStreamCreate(cudaStream_t* pStream)函数；
    1. 当执行一次异步数据传输时，我们必须(此处存疑)使用页锁定内存pinned memory或page-locked memory)：
    • 对CUDA架构而言，主机端的内存被分为两种，一种是可分页内存（pageable memroy）和页锁定内存（page-locked或pinned）。可分页内存是由操作系统API函数malloc()在主机内存上分配的，页锁定内存是由CUDA函数cudaHostAlloc()在主机内存上分配的。页锁定内存的重要属性是主机的操作系统将不会对这块内存进行分页（分页机制）和交换操作，确保该内存始终驻留在物理内存中。
    • 如果在异步数据传输期间，host端的某些操作导致内存换页，可能会导致数据传输异常（另外一个说法：没有显示定义页锁定内存时，当将pageable host Memory数据送到device时，CUDA驱动会分配临时的页锁定内存，并将host数据放到这个临时空间里，然后再向device传输数据）；
    • GPU知道页锁定内存的物理地址，可以通过直接内存访问\color{Blue}直接内存访问Direct Memory Access, DMA技术直接在主机和GPU之间复制数据，速率更快。由于每个页锁定内存都需要分配物理内存，并且这些内存不能交换到磁盘上，所以页锁定内存比使用标准malloc()分配的可分页内存更消耗内存空间。
    • 分配页锁定内存可以使用cudaHostAlloc()函数或cudaMallocHost()函数（C API）；
    • 释放页锁定内存使用cudaFreeHost()；
    • 将可分页内存注册为页锁定内存cudaHostRegister()；
    1. 关于多流，后续会做相关的测试，另外可以参考以下博客：

同步

  • cudaDeviceSynchronize()：该方法将阻塞CPU端线程的执行，直到GPU端完成之前CUDA的任务，包括kernel函数、数据拷贝等；
  • cudaStreamSynchronize()：这个方法接受一个stream ID，它将阻塞CPU执行直到GPU端完成相应stream ID的所有CUDA任务，但其它stream中的CUDA任务可能执行完也可能没有执行完；
  • cudaStreamQuery()：检查stream中的操作是否全部完成并返回状态，即使有操作没完成也不会阻塞host。如果所有操作都完成了，则返回cudaSuccess，否则返回cudaErrorNotReady；

事件（Event）

用来标记stream执行过程的某个特定的点，其主要用途是：1. 同步stream执行；2. 操控device运行步调。

    1. __host____device__cudaError_t cudaEventRecord ( cudaEvent_t event, cudaStream_t stream = 0 )：记录一个事件。如果stream是非零的，当流中所有的操作完毕，事件被记录；否则，当CUDA context中所有的操作完毕，事件被记录。由于这个操作是异步的，必须使用cudaEventQuery和/或cudaEventSyncronize函数来决定何时事件被真的记录了。如果cudaEventRecord 之前被调用了，并且事件还没有被记录，函数返回cudaErrorInvalidValue。
    1. cudaEventSyncronize()：阻塞host直到事件真的被记录。
    1. cudaError_t cudaEventElapsedTime(float* ms, cudaEvent_t start, cudaEvent_t stop)：返回start和stop之间的时间间隔，单位是毫秒。start和stop不必关联到同一个stream上，但是要注意，如果二者任意一个关联到了非默认stream上，时间间隔可能要比期望的大。这是因为cudaEventRecord是异步发生的，我们没办法保证度量出来的时间恰好就是两个event之间，所以只是想要gpu工作的时间间隔，则stop和strat都关联到默认stream就好了。
// create two events
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
// record start event on the default stream
cudaEventRecord(start);
// execute kernel
kernel<<<grid, block>>>(arguments);
// record stop event on the default stream
cudaEventRecord(stop);
// wait until the stop event completes
cudaEventSynchronize(stop);
// calculate the elapsed time between two events
float time;
cudaEventElapsedTime(&time, start, stop);
// clean up the two events
cudaEventDestroy(start);
cudaEventDestroy(stop);

参考