CUDA __threadfence()

kar*_*kar 18 cuda

我已经浏览了很多论坛帖子和NVIDIA文档,但我无法理解它是什么__threadfence()以及如何使用它.有人可以解释一下内在的目的是什么吗?

Cyg*_*sX1 45

通常,不能保证如果一个块将某些内容写入全局内存,则另一个块将"看到"它.除了发出它的块之外,对全局存储器的写入顺序也不保证.

有两个例外:

  • 原子操作 - 其他块始终可见
  • threadfence

想象一下,一个块产生一些数据,然后使用原子操作来标记数据存在的标志.但是有可能,另一个块将看到该标志,但会读取不正确或不完整的数据.

__threadfence函数停止当前线程,直到其对全局内存的写入保证被网格中的所有其他线程可见.所以,如果你做的事情如下:

  1. 存储您的数据
  2. __threadfence()
  3. 原子地标记一面旗帜

保证如果另一个块看到标志,它也会看到数据.

进一步阅读:Cuda编程指南,第B.2.4章和第B.5章

  • `__syncthreads()`强于`__threadfence_block()`.在__syncthreads()之后,您可以保证屏障之前所有共享/全局内存写入屏障之后的所有线程都可见.但是,`__syncthreads()`仅对块有影响,并且不在不同块的线程之间给出保证. (2认同)