2009-10-29 85 views
12

幾乎在任何地方,我讀過關於使用CUDA編程的內容,都提到重要性,即變形中的所有線程都執行相同的操作。
在我的代碼中,我有一種情況,我無法避免某種情況。它看起來像這樣:CUDA:同步線程

// some math code, calculating d1, d2 
if (d1 < 0.5) 
{ 
    buffer[x1] += 1; // buffer is in the global memory 
} 
if (d2 < 0.5) 
{ 
    buffer[x2] += 1; 
} 
// some more math code. 

一些線程可能會進入一個規定的條件,有些人可能會進入到這兩個和其他可能無法進入任一。

現在爲了讓所有線程在條件結束後再次回到「做同樣的事情」,我應該在使用__syncthreads()之後同步它們嗎?或者這是以某種方式自動發生的?
可以兩個線程是而不是做同樣的事情,由於其中一個是後面的一個操作,從而毀滅了每個人?還是有一些幕後的努力讓他們在分支之後再次做同樣的事情?

回答

35

在變形中,沒有線程會「超前」任何其他線程。如果有一個條件分支,並且它是由warp中的某些線程獲取的,但不是其他線程(又稱warp「divergence」),則其他線程將空閒直到分支完成,並且它們全部「聚合」回到一個共同指令。所以如果你只需要線程內部的同步,那就會「自動地」發生。

但是不同的經紗不是以這種方式同步的。因此,如果您的算法要求某些操作在許多經線上完成,那麼您需要使用顯式同步調用(請參見CUDA編程指南,第5.4節)。


編輯:重組未來數段,以澄清一些事情。

這裏真的有兩個不同的問題:指令同步和內存可見性。

  • __syncthreads()強制指令同步,確保存儲器的可見性,但只能在一個塊,而不是跨越塊(CUDA編程指南附錄B.6)。這對於在共享內存上執行寫入然後讀取很有用,但不適用於同步全局內存訪問。

  • __threadfence()確保全局內存可見性但不做任何指令同步,因此根據我的經驗,它的使用有限(但請參閱附錄B.5中的示例代碼)。

  • 全局指令同步是不可能的內核中。如果您在任何線程調用g()之前需要f()上的所有線程完成,拆分f()g()成兩個不同的內核,並從主機順序調用它們。

  • 如果您只需增加共享或全局計數器,請考慮使用原子增量函數atomicInc()(附錄B.10)。對於上面的代碼,如果x1x2不是全局唯一的(跨網格中的所有線程),非原子增量將導致競爭條件,類似於附錄B.2.4的最後一段。

最後,請記住,對全局存儲器和特別是同步功能(包括原子)的任何操作都不利於性能。

不知道你解決它的問題是很難猜測,但也許你可以重新設計你的算法使用共享內存而非全局內存在一些地方。這將減少同步的需要併爲您提供性能提升。

+0

讓我看看我是否明白。所以如果條件分支做相同數量的工作,這應該不會影響性能,因爲每個線程都不會空閒太久。我對嗎? – 2011-03-31 22:51:03

+1

@omegatai我認識到你的評論是舊的,但其他人可能想知道,所以在這裏:一個warp一次只能處理一條指令,所以如果warp中的某些線程做了一件事,其餘的線程做其他事情,總時間是這兩組線程的總時間量。在經紗內沒有時間重疊。性能受到影響。 – 2013-10-15 21:11:31

2

從CUDA最佳實踐指南的第6.1節:

任何流控制指令(如果,開關,做,對,而)可以顯著通過使相同的經紗的紗線以影響 的指令吞吐量發散;即 遵循不同的執行路徑。如果發生這種情況,必須對不同的執行路徑 進行串行化處理,從而增加爲此 warp執行的指令總數。當所有不同的執行路徑都完成後,線程會聚到 回到相同的執行路徑。

所以,你不需要做任何特別的事情。

1

您的問題的答案是否定的。你不需要做任何特別的事情。 無論如何,你可以解決這個問題,而不是你的代碼,你可以做這樣的事情:

buffer[x1] += (d1 < 0.5); 
buffer[x2] += (d2 < 0.5); 

你應該檢查是否可以在聚結的圖案使用共享內存和訪問全局內存。另外請確保您不想在多個線程中寫入相同的索引。

+0

訣竅很微妙,但通過這樣做,你根本不會分支! – purpletentacle 2014-11-14 16:13:04

2

在加布裏埃爾的迴應:

「環球指令同步是不可能的內核中如果需要調用任何線程克(),分割F()和g之前所有線程完成F()()分成兩個不同的內核,並從主機串行調用它們。「

如果你需要什麼?在同一個線程F()和g()的原因是因爲你使用的寄存器內存,你想從F寄存器或共享數據讓到g? 也就是說,對於我的問題,跨塊同步的全部原因是因爲f中的數據需要在g中 - 而向內核發展則需要大量額外的全局內存將寄存器數據從f傳輸到g,我想避免