等效于 CUDA 中的屏障（CLK_GLOBAL_MEM_FENCE）

Posted 2023-03-17

【中文标题】等效于 CUDA 中的屏障（CLK_GLOBAL_MEM_FENCE）

【英文标题】：Equivalent of barrier(CLK_GLOBAL_MEM_FENCE) in CUDA

【发布时间】：2019-01-23 15:45:14

【问题描述】：

在 CUDA 中调用 barrier(CLK_GLOBAL_MEM_FENCE) (OpenCL) 的等价物是什么？

它应该等到块中的所有线程都到达屏障。并且在屏障之前完成的全局内存访问应该对屏障之后块中的所有线程可见。

__syncthreads() 是否足够，或者全局内存栅栏需要__threadfence()，还是两者兼而有之？如果是这样，应该按什么顺序调用？

【参考方案1】：

您可能希望阅读有关内存栅栏和执行障碍的 CUDA 文档 here。

__syncthreads() 既是执行屏障（对于块中的线程），也是共享和全局内存操作的内存栅栏。对于全局内存操作，围栏强制仅针对块中的线程。

这里的栅栏有特定的含义： 在栅栏之前发生的内存操作将对块中的其他线程具有可见性，因为在栅栏之后发生的操作之前实际发生（即可见）。请查看已链接的文档以获取详细说明。

如果您还需要块外的（全局内存）防护（即网格中的所有线程），那么您还必须使用__threadfence()。与__syncthreads() 不同，__threadfence() 不是任何形式的执行障碍。它只是一个记忆栅栏。

调用顺序无关紧要（只要没有干预操作）。栅栏是一个分界点，__threadfence() 的功能是 __syncthreads() 中包含的栅栏功能的超集