GPU在用于通用计算时，非常强调与SIMD和SIMT的细粒度并行性。它们在具有高算术强度的常规数字处理工作负载上表现最佳。
尽管如此，为了适用于尽可能多的工作负载，它们还必须能够实现粗粒度 MIMD 并行性，其中不同的内核在不同的数据块上执行不同的指令流。
这意味着 GPU 上的不同内核在执行不同的指令流后必须相互同步。他们是怎么做到的？
在 CPU 上，答案是存在缓存一致性以及一组选择与 CAS 或 LL/SC 等良好配合使用的通信原语。但据我了解，GPU 没有缓存一致性 - 避免这样的开销是它们首先比 CPU 更高效的最大原因。
那么GPU核心之间用什么方法进行同步呢？如果它们如何交换数据的答案是通过写入共享主内存，那么它们如何同步，以便发送方可以通知接收方何时读取数据？
如果答案取决于特定的架构，那么我对支持 CUDA 的现代 Nvidia GPU 特别感兴趣。
编辑：从Booo链接的文档来看，这是我到目前为止的理解：
他们似乎使用"流"这个词来表示同步完成的大量工作(包括像SIMD这样的细粒度并行性);问题是如何在多个流之间同步/通信。
正如我推测的那样，这比在 CPU 上要明确得多。 特别是，他们谈到：
- 页面锁定内存
- cudaDeviceSyncing ()
- cudaStreamSync ( streamid )
- cudaEventSync (事件)
因此，流可以通过将数据写入主内存(或L3缓存？)进行通信，并且没有像CPU上的缓存一致性那样，而是有锁定的内存页和/或显式同步API。
我的理解是，有几种方法可以使用 CUDA "同步"：
-
CUDA 流(在功能级别)：
cudaDeviceSynchronize()在整个设备上同步。此外，您可以将特定流与cudaStreamSynchronize(cudaStream_t stream)同步，或者将嵌入在某些流中的事件与cudaEventSynchronize(cudaEvent_t event)同步。参考文献 1， 参考文献 2. -
协作组(>CUDA 9.0 和>CC 3.0)：您可以在组级别进行同步，组可以是一组合并线程、线程块或跨多个设备的网格。这要灵活得多。使用 定义您自己的组
(1) 当前合并的线程集的
auto group = cooperative_groups::coalesced_threads()，或(2)
auto group = cooperative_groups::this_thread_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()。参考文献 1， 参考文献 2.