2012-08-07 45 views
1

哪一個更好,單個Warp的線程之間或不同Warps線程之間的原子競爭(併發)?我認爲,當你訪問共享內存的時候,一個warp的線程彼此競爭的時候會比不同warp的線程少。而相反地訪問全局內存,最好是一個塊的不同線程的線程比單個線程的線程競爭更少,不是嗎?哪個更好,原子之間的競爭:單個Warp的線程還是不同的Warps線程?

我需要它來知道如何更好地解決競爭(併發)和更好地分隔存儲:單線程之間或線程之間的線程。

順便說一下,可以說團隊__ syncthreads();同步它在一個塊中變形而不是一個變形的線程?

回答

2

如果塊中有相當數量的線程對同一個值執行原子更新,那麼性能會很差,因爲這些線程都必須被序列化。在這種情況下,讓每個線程將其結果寫入單獨的位置,然後在單獨的內核中處理這些值通常會更好。

如果warp中的每個線程都執行原子更新爲相同的值,則warp中的所有線程都會在相同的時鐘週期內執行更新,因此它們必須在原子更新點處全部序列化。這可能意味着該warp預定了32次以獲得所有服務的線程(非常糟糕)。另一方面,如果塊中的每個變形中的單個線程執行原子更新爲相同的值,則該影響將會較低,這是因爲經線對(在兩個經線在每個時鐘處處理的兩個經線調度程序)在時間上偏移(一個時鐘週期),因爲它們在處理管道中移動。所以最終只有兩個原子更新(兩個經線中的每一個),在一個週期內發佈並需要立即被序列化。

因此,在第二種情況下,情況較好,但仍存在問題。原因在於,取決於共享值的位置,您仍然可以在SM之間進行序列化,並且這可能非常緩慢,因爲每個線程可能必須等待更新才能全部移出到全局內存,或者至少L2,然後回來。有可能以這樣的方式重構算法,即塊內的線程對共享內存(L1)中的值執行原子更新,然後每個塊中的一個線程對全局內存中的值執行原子更新(L2 )。

原子操作可以是完整的救生員,但他們往往被CUDA新手過度使用。通常使用單獨步驟並行簡化或並行流壓縮算法更好(請參閱thrust::copy_if)。

相關問題