CUDA中的流的相关要点


总结一下流的相关要点。

1、流中排队的所有操作,相对于主机都是异步的,但在流的内部,它们严格按照主机代码确定的顺序执行的。

2、异步函数和流是在cuda中构建网格级并发的两个基本支柱。

3、cuda中的所有操作(包括内核和数据传输)都在一个流中显示或隐式地运行。

4、流分默认和非默认两种,也叫空流和非空流,还叫隐式声明的流和显示声明的流,都是一个意思,按照默认和非默认来理解最直观。

5、在使用异步cuda函数时,常见的疑惑在于,它们可能会从同一个流中在它们前面启动的异步操作中返回错误代码。因此返回错误的API调用并不一定是产生错误的那个调用。

6、流的基本操作

cudaStream_t stream;
cudaStreamCreate(&stream);
kernel_name<<<grid, block, sharedMemSize, stream>>>(argument list);

cudaError_t cudaStreamSynchronize(cudaStream_t stream);
cudaError_t cudaStreamQuery(cudaStream_t stream);

7、影响并发内核最大数量的因素主要是设备的硬件工作队列的数量,Hyper-Q技术使用多个硬件工作队列,从而避免了单个硬件工作队列对多个流并发的限制(即虚假依赖)。Hyper-Q技术通过在主机和设备之间维持多个硬件管理上的连接,允许多个cpu线程或进程在单一gpu上同时启动工作。

8、cuda中的事件本质上是对cuda流中的标记,它与该流内操作中的特定点相关联。事件的基本作用有两个:

  • 同步流的执行
  • 监控设备的进展

cuda的api提供了在流中任意点插入事件以及查询事件完成的函数。

事件操作的基本代码

cudaEvent_t event;
cudaError_t cudaEventCreate(cudaEvent_t* event);
cudaError_t cudaEventDestroy(cudaEvent_t event);
cudaError_t cudaEventRecord(cudaEvent_t event, cudaStream_t stream = 0);

cudaError_t cudaEventSynchronize(cudaEvent_t event);
cudaError_t cudaEventQuery(cudaEvent_t event);

cudaError_t cudaEventElapsedTime(float* ms, cudaEvent_t start, cudaEvent_t stop);

事件的启动和停止不必在同一个cuda流中。

9、非默认流分为两种,阻塞流和非阻塞流。对于阻塞流,在主机代码中添加到默认流的操作,会把后面加入阻塞流的操作阻塞。非阻塞流不会受限制。默认的流是阻塞流。

创建阻塞和非阻塞流的代码:

cudaError_t cudaStreamCreateWithFlags(cudaStream_t* pStream, unsigned int flags);

//flags取值为:
//cudaStreamDefault
//cudaStreamNonBlocking

10、隐式同步的常见操作:

  • 锁页主机内存分配
  • 设备内存分配
  • 设备内存初始化
  • 同一设备上两个地址之间的内存复制
  • 一级缓存/共享内存配置的修改

11、显式同步的常见操作:

  • 同步设备
  • 同步流
  • 同步流中的事件
  • 使用事件跨流同步

流间同步代码:

cudaError_t cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event);

该函数能使指定流等待指定事件。该事件可能与同一个流相关,也可能与不同的流相关。

12、使用OpenMP的调度操作 OpenMP是CPU的并行编程模型,它使用编译器指令来识别并行区域。在使用OpenMP的同时使用cuda,不仅可以提高便携性和生产效率,而且还可以提高主机代码的性能。 如果每个流在内核执行前,期间或之后有额外的工作待完成,那么它可以包含在同一个OpenMP并行区域内,并且跨流和线程进行重叠。

13、重叠内核执行和数据传输 数据传输和内核执行有两种关系:

  • 如果内核使用数据A,那么对A进行数据传输必须要安排在内核启动前,且必须位于相同的流中。
  • 如果一个内核完全不使用数据A,那么内核执行和数据传输可以位于不同的流中。

14、流回调是一种可以到cuda流中排列等待的操作。一旦流回调之前所有的流操作全部完成,被流回调指定的主机端函数就会被cuda运行时所调用,回调函数最终还是在主机上执行。

添加回调的代码:

cudaError_t cudaStreamAddCallback(cudaStream_t stream, cudaStreamCallback_t callback, void *userData, unsigned int flags);

每使用cudaStreamAddCallback一次,只执行一次回调,并阻塞队列中排在后面的工作,直到回调函数完成。

对回调函数有两个限制:

  • 从回调函数中不可以调用cuda的api函数
  • 在回调函数中不可以执行同步

15、在cuda中,通常可以使用3种不同类型的重叠方案来隐藏计算或通信延迟。

  • 在设备上重叠多个并发的内核
  • 重叠带有传入或传出设备数据传输的cuda内核
  • 重叠cpu执行和gpu执行

Refs:

《cuda c编程 权威指南》


原创文章,转载请注明出处,否则拒绝转载!
本文链接:抬头看浏览器地址栏

上篇: AI经典资料收集
下篇: CUDA多GPU编程要点