CUDA中的threadfence内置函数的目的是什么？

Question

我已经阅读了许多论坛帖子和NVIDIA文档，但是我无法理解__threadfence()的作用以及如何使用它。有人能解释一下这个内置函数的目的吗？

- kar

Answer 1

通常情况下，如果一个块向全局内存中写入了某些内容，另一个块并不能保证“看到”它。此外，在写入全局内存时也没有任何保证，除了发出该写入的块之外。

有两个例外：

- 原子操作 - 这些始终可以被其他块看到 - threadfence 想象一下，一个块生成了一些数据，然后使用原子操作标记一个标志，表示数据已经存在。但是，其他块在看到标志后仍可能读取不正确或不完整的数据。

__threadfence函数来拯救，确保顺序。在它之前的所有写入都真正发生在它之前，在其他块中看到。

请注意，__threadfence函数不一定需要阻塞当前线程，直到它对全局内存的写入对网格中的所有其他线程可见为止。以这种天真的方式实现，__threadfence函数可能会严重影响性能。

例如，如果您执行以下操作：

1. 存储数据 2. __threadfence() 3. 原子标记一个标志

则保证如果其他块看到标志，则它也将看到数据。

进一步阅读：CUDA编程指南，第B.5章（截至版本11.5）

- CygnusX1

---

__syncthreads()比__threadfence_block()更强。在__syncthreads()之后，您可以保证屏障之前的所有共享/全局内存写入对屏障之后的所有线程都是可见的。但是，__syncthreads()仅对块产生影响，并且不保证不同块的线程之间的任何内容。 - CygnusX1

其他SM的L1会与存储的数据协调吗？还是您仍需要指定全局作用域加载（L1-非缓存）？ - maxbc

@user2023370 不是的。__syncthreads 和 __threadfence_block 确保给定块中所有先前的全局更改对该块中的所有线程可见。然而，这通常不是 syncthreads 的常见用例，因为通常你会想要使用共享内存来进行这种类型的线程通信。 - CygnusX1

B.2.4现在是：https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#managed - Jigao Luo

@LuoJigao更新了B.5章节并添加了一个链接。我不记得10年前我写的B.2.4是什么，为了避免混淆，我将其删除了。 - CygnusX1