【知识】CUDA Stream和并发性

转载请注明出处：小锋学长生活大爆炸[xfxuezhagn.cn] 如果本文帮助到了你，欢迎[点赞、收藏、关注]哦~

翻译来自： Streams And Concurrency Webinar

https://developer.download.nvidia.com/CUDA/training/StreamsAndConcurrencyWebinar.pdf 可以看原版pdf，本博客只是挑选并翻译一下。 更详细的，推荐看这个编程指南：https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#streams

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#streams

支持同时执行多个CUDA操作的能力 (超越了多线程并行性)

• CUDA Kernel <<<>>>
• cudaMemcpyAsync (HostToDevice)
• cudaMemcpyAsync (DeviceToHost)
• CPU上的操作

Fermi架构可以同时支持 (要求计算能力2.0+)

• 在GPU上最多有16个CUDA Kernel
• 2个cudaMemcpyAsyncs(必须是在不同的方向上)
• CPU上的计算

Stream：在GPU上按问题顺序执行的一系列操作。用于影响并发性的编程模型：

• 不同流中的CUDA操作可以并发地进行
• 来自不同流的CUDA操作可以交错进行

默认Stream (又称为 Stream 0)，即未指定流时使用的流。

并发性要求

• CUDA操作必须在不同的、非0的流中；
• cudaMemcpyAsync时是主机中的'pinned' memory；
• 必须有足够的资源：
  • cudaMemcpyAsyncs在不同的方向上；
  • 设备资源(SMEM、registers、blocks等）；

简单的同步示例：默认流中的所有CUDA操作都是同步的

简单的异步且没有流的示例：默认情况下，GPU内核与主机是异步的

简单的使用了流的异步示例：完全异步/并发，并发操作所使用的数据需要是独立的

显式同步的情况

• 同步所有内容
  • cudaDeviceSynchronize ()
  • 阻止主机，直到所有发出的CUDA调用都完成
• 特定的流中同步w.r.t.
  • cudaStreamSynchronize ( streamid )
  • 阻塞主机，直到流中的所有CUDA调用完成
• 使用事件同步
  • 在流中创建特定的“事件”，以用于同步
  • cudaEventRecord ( event, streamid )
  • cudaEventSynchronize ( event )
  • cudaStreamWaitEvent ( stream, event )
  • cudaEventQuery ( event )

显式同步示例：使用事件解决

隐式同步：这些操作会隐式地同步所有其他CUDA操作

• 锁定页面的内存分配
  • cudaMallocHost
  • cudaHostAlloc
• 设备内存分配
  • cudaMalloc
• 内存操作的非异步版本
  • cudaMemcpy* (no Async suffix)
  • cudaMemset* (no Async suffix)
• 更改为L1/共享内存配置
  • cudaDeviceSetCacheConfig

流的调度

• Fermi硬件有3个队列
  • 1个计算引擎队列
  • 2个复制引擎队列：一个用于H2D，一个用于D2H
• CUDA操作按它们被发布的顺序被发送到HW
  • 被放置在相关的队列中
  • 维护了引擎队列之间的流依赖关系，但在引擎队列中丢失
• 如果出现以下情况，则将从引擎队列分派CUDA操作：
  • 在同一流中的先前的调用已经完成
  • 在同一流中的先前的调用已经分派
  • 资源可用
• 如果CUDA内核在不同的流中，则可以并发地执行它们
  • 如果前面内核的所有线程块都已调度，并且仍然有SM资源可用，则调度给定内核的线程块
• 请注意，一个被阻塞的操作会阻止队列中的所有其他操作，甚至在其他流中

Blocked Queue的示例

最优并发性可能取决于内核的执行时间

并发内核的调度

• 并发内核调度很特殊
• 通常，在操作后将信号插入队列，在同一流中启动下一个操作
• 对于计算引擎队列，若要启用并发内核，当计算内核按顺序发出时，该信号将延迟到最后一个顺序计算内核之后
• 在某些情况下，这种信号的延迟可能会阻塞其他队列

并发内核和阻塞的示例

其他的详细信息

• 很难同时运行超过4个内核
• 可以使用环境变量禁用并发性：CUDA_LAUNCH_BLOCKING
• cudaStreamQuery可以用来分离连续的核，防止延迟信号
• 使用超过8个textures的内核不能并发运行
• 切换L1/共享配置将破坏并发性
• 要并发运行，CUDA操作必须有不超过62个介入的CUDA操作。也就是说，按照“问题顺序”，它们不能被超过62个其他问题分开，进一步的操作会被序列化。
• cudaEvent_t对于计时很有用，但对于性能的使用可以用cudaEventCreateWithFlags ( &event, cudaEventDisableTiming )