OpenCL work_group_barrier()的CUDA等效实现：支持超工作组范围？

问题背景 #

《OpenCL 3.0 C规范》中的内置函数work_group_barrier(cl_mem_fence_flags flags, memory_scope scope)（本次忽略flags参数），可确保工作组内所有工作项的内存操作对指定scope可见。可选的scope包括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()可满足前三种scope的要求，请问针对后三种（设备级、所有SVM设备、所有设备）范围，是否存在具备对应或更强语义的CUDA内置屏障函数？

对应解决方案 #

1. 设备级（memory_scope_device） #

CUDA 9.0及以上提供协作组（Cooperative Groups） API，其中this_grid().sync()可实现当前网格内所有线程的同步，等价于设备级屏障语义（前提是网格内所有线程归属同一设备）。此外，若仅需保证内存操作对全设备可见而非强制线程同步，可使用内核函数__threadfence()，它能确保当前线程的内存写操作对设备内所有其他线程可见。

2. 所有SVM设备（memory_scope_all_svm_devices） #

CUDA没有单一内置核函数直接实现该范围的屏障，需通过多步骤组合实现：

• 先通过协作组this_grid().sync()或主机端cudaDeviceSynchronize()完成单设备内的同步；
• 结合**共享虚拟内存（SVM）**或统一内存的一致性机制，配合主机端cudaMemAdvise()等操作，保证跨SVM设备的内存操作可见性；
• 最后调用cudaDeviceSynchronize()同步所有参与的SVM设备。

3. 所有设备（memory_scope_all_devices） #

该范围涉及跨所有GPU设备甚至CPU的全局同步，CUDA无内核内单一函数可直接实现。需通过主机端协调：

• 对每个设备调用cudaDeviceSynchronize()，确保单设备内所有内核执行完成；
• 依托统一内存或SVM的跨设备内存一致性机制，保证内存操作在所有设备间可见；
• 若需严格的线程级跨设备同步，需搭配主机端同步原语（如互斥锁）完成协调。

内容的提问来源于stack exchange，提问作者einpoklum