【发布时间】：2011-07-11 02:25:22
【问题描述】：
我浏览了许多论坛帖子和 NVIDIA 文档，但我无法理解 __threadfence() 的作用以及如何使用它。有人能解释一下这个内在的目的是什么吗？
【问题讨论】：
标签： cuda
【发布时间】：2011-07-11 02:25:22
【问题描述】：
通常，不能保证如果一个块将某些内容写入全局内存，另一个块会“看到”它。也不能保证写入全局内存的顺序，发出它的块除外。
有两个例外：
- 原子操作 - 其他方块始终可见
- 线程栅栏
想象一下，一个块产生一些数据，然后使用原子操作标记数据存在的标志。但是有可能其他block在看到flag之后仍然读取到不正确或者不完整的数据。
__threadfence 函数来救援，确保排序。从其他块可以看出，它之前的所有写入确实发生在它之后的所有写入之前。
请注意，__threadfence 函数不一定需要暂停当前线程，直到它对全局内存的写入对网格中的所有其他线程都是可见的。以这种幼稚的方式实现，__threadfence 函数可能会严重影响性能。
例如，如果您执行以下操作：
- 存储您的数据
__threadfence()- 原子标记标志
保证如果其他块看到标志，它也会看到数据。
延伸阅读：Cuda 编程指南，第 B.5 章（截至 11.5 版）
【讨论】：
-
__syncthreads() 怎么样，是否保证块中任何线程的内存访问对块中的所有线程都是可见的？
-
__syncthreads()比__threadfence_block()强。在__syncthreads()之后，您可以保证屏障之前的所有共享/全局内存写入对屏障之后的所有线程都是可见的。但是，__syncthreads()仅对块有影响，并且不保证不同块的线程之间存在。 -
@user2023370 No.
__syncthreads和__threadfence_block确保该块中的所有线程都可以看到给定块的所有先前全局更改。然而，这对于同步线程来说不太常见，因为通常您希望将共享内存用于这种线程通信。 -
@LuoJigao 更新了 B.5 章以包含链接。我不记得 10 年前写的时候 B.2.4 是什么，为了不引起混淆，我删除了它。