2011-03-08 62 views
15

我經歷了很多論壇和NVIDIA手冊,但我不明白什麼是 __threadfence()和它的使用?CUDA __threadfence()

謝謝。

回答

40

通常情況下,如果一個塊向全局內存寫入某些內容,則不能保證其他塊會「看到」它。寫入全局內存的順序也不能保證,除了發佈它的塊之外。

有兩個例外:

  • 原子操作 - 這些都是始終可見通過其它塊
  • threadfence

可以想象,一個塊產生一些數據,然後使用原子操作來標記數據在那裏的標誌。但有可能,另一個塊會看到該標誌,但會讀取不正確或不完整的數據。

__threadfence函數會暫停當前線程,直到它寫入全局內存時保證它可以被網格中的所有其他線程看到。所以,如果你這樣做:

  1. 存儲你的數據
  2. __threadfence()
  3. 原子標記標誌

可以保證的是,如果其他塊看到標誌,它也將看到數據。

延伸閱讀:CUDA編程指南,章節B.2.4和B.5

+0

如何__syncthreads(),是否保證內存塊中的任何線程訪問將是可見的塊中的所有線程? – 2013-05-14 21:11:19

+2

'__syncthreads()'強於'__threadfence_block()'。在'__syncthreads()'之後,你可以保證所有的共享/全局內存在屏障之前的所有線程都可見之前寫入。然而'__syncthreads()'只對一個塊有影響,並且在不同塊的線程之間沒有保證。 – CygnusX1 2013-05-14 22:12:52

+0

來自其他SM的L1將與存儲的數據保持一致嗎?或者你是否仍然需要指定全局範圍加載(L1-non-cacheable)? – maxbc 2017-08-28 23:07:37