OpenCL 3.0 C规范 包含一个内在函数

work_group_barrier(
    cl_mem_fence_flags  flags,
    memory_scope        scope
);

在本题的目的下，我们先忽略 flags 参数。

当一个工作项在调用该函数后继续执行时，可以确保对所指定范围内的任意内存操作（它自己的以及同一工作组中其他工作项的内存操作）都是可见的。可能的作用域如下：

memory_scope_work_item
memory_scope_sub_group
memory_scope_work_group
memory_scope_device
memory_scope_all_svm_devices
memory_scope_all_devices

现在，在一个CUDA内核中调用 __syncthreads()，这将满足（若我没理解错的话）前3 个作用域的要求。但是，最后3 个作用域呢？ 整个设备、所有共享虚拟内存设备，以及所有设备？是否存在具有这些（或更强）语义的CUDA内置屏障吗？

解决方案

用于整个设备的内存屏障类型是 cuda::thread_scope_device，在此之上还有 cuda::thread_scope_system；其描述为

内存操作对同一系统中的其他线程可见（CPU、其他GPU）。

__syncthreads() 屏障本身并没有那么强。它只为参与的线程（块级）提供一个内存屏障。它的描述 仅提到此级别的内存排序。

因此你需要通过 thread fence via cuda::atomic_thread_fence(cuda::memory_order_..., cuda::thread_scope_system) 或 __threadfence_system() 再添加一个额外的屏障。

然而，我不确定把这与 __syncthreads() 搭配使用一定是正确的做法。你确定需要一个块级执行屏障 再加上 一个系统级内存屏障吗？也许你更适合使用其他同步原语之一，例如屏障、信号量或简单原子变量等，它们都可以与系统级作用域一起工作。