GPU 核心如何相互通信？
How do GPU cores communicate with each other?
GPU 在用于通用计算时，非常重视与 SIMD 和 SIMT 的细粒度并行性。它们在具有高算术强度的常规数字运算工作负载上表现最佳。
尽管如此，要适用于尽可能多的工作负载，它们还必须能够进行粗粒度 MIMD 并行处理，其中不同的内核在不同的数据块上执行不同的指令流。
这意味着GPU上的不同核心在执行不同的指令流后必须相互同步。他们是怎么做到的？
在 CPU 上，答案是存在缓存一致性以及一组通信原语，这些原语被选择为与 CAS 或 LL/SC 等一起工作。但据我所知，GPU 没有缓存一致性——避免这种开销是它们比 CPUs 更高效的最大原因。
那么GPU核之间是通过什么方式进行同步的呢？如果他们如何交换数据的答案是写入共享主内存，那么他们如何同步以便发送方可以通知接收方何时读取数据？
如果答案取决于特定架构，那么我对支持 CUDA 的现代 Nvidia GPU 特别感兴趣。
编辑：来自链接的文档 Booo，这是我目前的理解：
他们似乎使用 'stream' 这个词来表示大量同步完成的事情（包括像 SIMD 这样的细粒度并行性）；那么问题是如何在多个流之间synchronize/communicate。
正如我推测的那样，这比 CPU 上的要明确得多。特别是，他们谈论：
- 页面锁定内存
- cudaDeviceSynchronize()
- cudaStreamSynchronize ( streamid )
- cudaEventSynchronize（事件）
所以流可以通过将数据写入主内存（或 L3 缓存？）来进行通信，并且没有什么比 CPUs 上的缓存一致性更好的了，取而代之的是锁定内存页，and/or 显式同步 API.
我的理解是，有几种方法可以使用 CUDA 进行“同步”：
CUDA 流（在功能级别）：cudaDeviceSynchronize() 在整个设备上同步。此外，您可以将特定流与 cudaStreamSynchronize(cudaStream_t stream) 同步，或将嵌入在某些流中的事件与 cudaEventSynchronize(cudaEvent_t event) 同步。 Ref 1, Ref 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() 在共享内存中执行传统的块级同步。 Ref 1, Ref 2.
GPU 在用于通用计算时，非常重视与 SIMD 和 SIMT 的细粒度并行性。它们在具有高算术强度的常规数字运算工作负载上表现最佳。
尽管如此，要适用于尽可能多的工作负载，它们还必须能够进行粗粒度 MIMD 并行处理，其中不同的内核在不同的数据块上执行不同的指令流。
这意味着GPU上的不同核心在执行不同的指令流后必须相互同步。他们是怎么做到的？
在 CPU 上，答案是存在缓存一致性以及一组通信原语，这些原语被选择为与 CAS 或 LL/SC 等一起工作。但据我所知，GPU 没有缓存一致性——避免这种开销是它们比 CPUs 更高效的最大原因。
那么GPU核之间是通过什么方式进行同步的呢？如果他们如何交换数据的答案是写入共享主内存，那么他们如何同步以便发送方可以通知接收方何时读取数据？
如果答案取决于特定架构，那么我对支持 CUDA 的现代 Nvidia GPU 特别感兴趣。
编辑：来自链接的文档 Booo，这是我目前的理解：
他们似乎使用 'stream' 这个词来表示大量同步完成的事情（包括像 SIMD 这样的细粒度并行性）；那么问题是如何在多个流之间synchronize/communicate。
正如我推测的那样，这比 CPU 上的要明确得多。特别是，他们谈论：
- 页面锁定内存
- cudaDeviceSynchronize()
- cudaStreamSynchronize ( streamid )
- cudaEventSynchronize（事件）
所以流可以通过将数据写入主内存（或 L3 缓存？）来进行通信，并且没有什么比 CPUs 上的缓存一致性更好的了，取而代之的是锁定内存页，and/or 显式同步 API.
我的理解是，有几种方法可以使用 CUDA 进行“同步”：
CUDA 流（在功能级别）：
cudaDeviceSynchronize()在整个设备上同步。此外，您可以将特定流与cudaStreamSynchronize(cudaStream_t stream)同步，或将嵌入在某些流中的事件与cudaEventSynchronize(cudaEvent_t event)同步。 Ref 1, Ref 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()在共享内存中执行传统的块级同步。 Ref 1, Ref 2.