如果塊中有相當數量的線程對同一個值執行原子更新,那麼性能會很差,因爲這些線程都必須被序列化。在這種情況下,讓每個線程將其結果寫入單獨的位置,然後在單獨的內核中處理這些值通常會更好。
如果warp中的每個線程都執行原子更新爲相同的值,則warp中的所有線程都會在相同的時鐘週期內執行更新,因此它們必須在原子更新點處全部序列化。這可能意味着該warp預定了32次以獲得所有服務的線程(非常糟糕)。另一方面,如果塊中的每個變形中的單個線程執行原子更新爲相同的值,則該影響將會較低,這是因爲經線對(在兩個經線在每個時鐘處處理的兩個經線調度程序)在時間上偏移(一個時鐘週期),因爲它們在處理管道中移動。所以最終只有兩個原子更新(兩個經線中的每一個),在一個週期內發佈並需要立即被序列化。
因此,在第二種情況下,情況較好,但仍存在問題。原因在於,取決於共享值的位置,您仍然可以在SM之間進行序列化,並且這可能非常緩慢,因爲每個線程可能必須等待更新才能全部移出到全局內存,或者至少L2,然後回來。有可能以這樣的方式重構算法,即塊內的線程對共享內存(L1)中的值執行原子更新,然後每個塊中的一個線程對全局內存中的值執行原子更新(L2 )。
原子操作可以是完整的救生員,但他們往往被CUDA新手過度使用。通常使用單獨步驟並行簡化或並行流壓縮算法更好(請參閱thrust::copy_if
)。