multithreading - GPU 核心如何相互通信？

Question

GPU 在用于通用计算时，非常强调与 SIMD 和 SIMT 的细粒度并行性。它们在具有高算术强度的常规数字运算工作负载上表现最佳。

尽管如此，要适用于尽可能多的工作负载，它们还必须能够进行粗粒度 MIMD 并行处理，其中不同的内核在不同的数据 block 上执行不同的指令流。

这意味着 GPU 上的不同内核在执行不同的指令流后必须相互同步。他们是怎么做到的？

在 CPU 上，答案是缓存一致性加上一组通信原语，这些原语被选择用于与 CAS 或 LL/SC 等通信原语很好地协同工作。但据我了解，GPU 没有缓存一致性 - 避免此类开销是它们比 CPU 更高效的最大原因。

那么 GPU 内核之间使用什么方法进行同步呢？如果他们如何交换数据的答案是通过写入共享主内存，那么他们如何同步以便发送方可以通知接收方何时读取数据？

如果答案取决于特定的架构，那么我对支持 CUDA 的现代 Nvidia GPU 特别感兴趣。

编辑:从链接的文档 Booo 中，这是我目前的理解:

他们似乎使用“流”这个词来表示大量同步完成的事情(包括像 SIMD 这样的细粒度并行性)；那么问题就是如何在多个流之间进行同步/通信。

正如我推测的那样，这比在 CPU 上更明确。特别是，他们谈论:

• 页锁内存
• cudaDeviceSynchronize()
• cudaStreamSynchronize ( streamid )
• cudaEventSynchronize(事件)

因此流可以通过将数据写入主内存(或 L3 缓存？)来进行通信，并且没有像 CPU 上的缓存一致性那样的东西，取而代之的是锁定内存页面和/或显式同步 API。

Answer 1

我的理解是有几种方法可以使用 CUDA 进行“同步”:

• CUDA 流(在函数级别):cudaDeviceSynchronize()在整个设备上同步。此外，您可以将特定流与 cudaStreamSynchronize(cudaStream_t stream) 同步, 或者将嵌入在某些流中的事件与 cudaEventSynchronize(cudaEvent_t event) 同步. Ref 1 , Ref 2 .

• 协作组(>CUDA 9.0 和 >CC 3.0):您可以在组级别进行同步，一个组可以是一组合并的线程、一个线程 block 或跨越多个设备的网格。这要灵活得多。使用

  定义你自己的组

  (1) auto group = cooperative_groups::coalesced_threads()对于当前合并的线程集，或者

  (2) auto group = cooperative_groups::this_thread_block()对于当前线程 block ，您可以在 block 内进一步定义细粒度组，例如 auto group_warp = cooperative_groups::tiled_partition<32>(group) , 或者

  (3) auto group = cooperative_groups::this_grid()或 auto group = cooperative_groups::this_multi_grid()用于跨多个设备的网格。

  然后，您可以调用 group.sync()用于同步。您需要有支持 cooperativeLaunch 的设备或 cooperativeMultiDeviceLaunch通过。请注意，对于协作组，您已经可以使用 __syncthreads() 在共享内存中执行传统的 block 级同步。 . Ref 1 , Ref 2 .