缺少__syncthreads（）是否會阻止讀寫後的warp級共享內存危害？

我有一個warp，它將一些數據寫入共享內存 - 沒有覆蓋，並且在從共享內存讀取後不久。雖然在我的區塊中可能會有其他的變形，但它們不會觸及該共享內存的任何部分，也不會寫入任何我感興趣的翹曲讀取的地方。缺少__syncthreads（）是否會阻止讀寫後的warp級共享內存危害？

現在，我還記得，儘管步調一致執行經線，我們不能保證共享內存讀取之後的共享內存寫入將返回早些時候經理應寫入各自的值。（理論上這可能是由於指令重新排序或者 - 如@ RobertCrovella指出的那樣 - 編譯器優化共享內存訪問）

所以，我們需要求助於某些顯式的同步。很顯然，塊級__syncthreads()工作。這就是does：

__syncthreads()用於協調同一個塊的線程之間的通信。當一個塊中的某些線程訪問共享或全局內存中的相同地址時，對於這些內存訪問中的一些，可能存在寫後讀，寫後讀或後寫後寫入危險。通過同步這些訪問之間的線程可以避免這些數據危害。

這太厲害了我的需求：

它適用於全球存儲器中，而不僅僅是共享內存。
它執行幀間經線同步;我只需要intra-warp。
它可以防止所有類型的危害R-後-W，W-後-R，W-後-W;我只需要R-after-W。
它的工作原理也可用於多線程執行的情況下，寫入到共享存儲器相同的位置;在我的情況下，所有共享內存寫入不相交。

另一方面，像__threadfence_block()does not seem to suffice。這兩個級別的力量之間有什麼「中間」？

注：

相關問題：CUDA __syncthreads() usage within a warp。
如果你要建議我使用洗牌代替，那麼，是的，這有時可能 - 但如果你想擁有的數據數組訪問，即動態地決定你要去的共享數據的元素讀。這可能會蔓延到當地的記憶，這對我來說似乎很可怕。
我想也許volatile可能對我非常有用，但我不知道，如果使用它會做我想做的。
如果你有一個答案，假設計算機的能力至少是XX.YY，那就夠用了。

來源

2017-04-20 einpoklum

對於經紗級使用，根據第一段中的描述，「volatile」可能是有意義的。沒有一個清晰的例子是不可能的。根據您的描述，我沒有看到這與[古典並行壓縮教程]（http://developer.download.nvidia.com/assets/cuda/files/reduction.pdf）中的使用方式有什麼不同）（幻燈片22）。底層問題不是指令重新排序，而是編譯器優化共享值到寄存器 –

@RobertCrovella：這樣的編譯器優化不會確保我不需要'__syncthreads（）'開始？實際上 - 當它知道其他線程從編譯時不確定的共享內存位置讀取時，它怎麼可能進行這種優化？ – einpoklum

我不明白第一個問題。關於第二個問題，編譯器不知道其他線程正在做什麼。 'volatile'是你與編譯器交流的方式，其他線程可能會觸及我正在看的這個東西。這是關於GPU編譯器的常見誤解 - 它以某種方式感知多線程，並且其行爲將根據線程間可能發生的情況而改變。它不這樣工作。編程模型有一個單一的線程，編譯器也是如此。 –

如果我理解正確@RobertCrovella，這個代碼片段應該是安全的危害：

/* ... */ 
volatile MyType* ptr = get_some_shared_mem(); 
ptr[lane::index()] = foo(); 
auto other_lane_value = ptr[bar() % warpSize]; 
/* ... */

因爲使用volatile的。

來源

2017-04-20 13:32:44 einpoklum

缺少__syncthreads（）是否會阻止讀寫後的warp級共享內存危害？

回答

相關問題