經紗如何與原子操作一起工作？

Question

經線中的線程在物理上並行運行，因此如果其中一個線程（稱爲線程X）啓動了一個原子操作，那麼其他人會做什麼？等待？是不是意味着所有的線程都會等待，而線程X被推入原子隊列，獲得訪問（mutex）並執行一些有內存的東西，這個東西被這個互斥鎖保護着，並且之後會釋放互斥鎖？

有什麼辦法可以爲某些工作採取其他線程，如讀取一些內存，所以原子操作會隱藏它的延遲？我的意思是，15個空閒線程是..不好，我猜。原子很慢，是嗎？我如何加速它？有什麼模式可以使用它嗎？

對於銀行或整個內存是否具有共享內存鎖的原子操作？ 例如（不mutexs），有__shared__ float smem[256];

• 線程1運行atomicAdd(smem, 1);
• 線程2運行atomicAdd(smem + 1, 1);

那些線程與不同的銀行工作，但在一般的共享內存。他們是否運行parralel或他們將排隊？如果Thread1和Thread2來自單獨的warps或一般的warp，這個例子有什麼不同嗎？

Answer 1

我計算了10個問題。這使得很難回答。建議你問每個問題一個問題。

一般而言，warp中的所有線程都執行相同的指令流。所以有兩種情況可以考慮：

1. 沒有條件語句（例如，如果......那麼......否則）在這種情況下，所有線程都執行相同的指令，這恰好是一個原子指令。然後，所有32個線程將執行一個原子，儘管不一定在同一個位置。所有這些原子將被SM處理，並在一定程度上將序列化（如果它們正在更新相同的位置，它們將完全序列化）。
2. 帶條件例如，假設我們有if (!threadIdx.x) AtomicAdd(*data, 1);然後，線程0將執行原子，而其他線程不會執行 。看起來好像我們可以讓別人做 別的事情，但鎖步執行不允許這樣做。 跳轉執行是序列化的，使得所有執行if (true)路徑的線程將一起執行，並且執行該路徑的所有線程將一起執行，但是真實和錯誤的路徑將被序列化。所以再次，我們不可能真的有不同的 線程在執行不同的指令同時執行。

它的網絡是，在一個warp內，我們不能讓一個線程做一個原子而另一個線程同時做別的事情。

許多其他問題似乎預計內存事務在它們起源的指令週期結束時完成。事實並非如此。對於全局和共享內存，我們必須在代碼中採取特殊步驟，以確保先前的寫入事務對其他線程可見（這可以被視爲交易完成的證據。要做到這一點），一個典型的方法是使用屏障指令，例如__syncthreads()或__threadfence()但是，如果沒有這些屏障指令，線程不被「等待」的寫入完成。 A（依賴於a的操作）讀取可能會阻塞一個線程。寫入通常不能拖延一個線程。

現在讓我們看看你的問題：

所以，如果他們中的一個開始一個原子操作，會做什麼其他的？等待？

不，他們不會等待。原子操作被分派到處理原子的SM上的一個功能單元，並且所有線程一起以鎖步方式繼續。由於原子通常意味着讀取，是的，讀取可以停止扭曲。但是線程不會等待原子操作完成（即寫入）。然而，這個位置的後續讀操作會搪塞經，再次等待原子（寫入）來完成。在保證更新全局內存的全局原子的情況下，它將使原始SM（如果啓用）和L2（如果它們包含該位置作爲條目）中的L1無效。

有什麼辦法可以讓其他線程進行一些工作，比如讀取一些內存，所以原子操作會隱藏它的延遲？

不是真的，理由我在開始時說。

原子實在是太慢了，不是嗎？我如何加速它？有什麼模式可以使用它嗎？

是的，原子公司可以使程序運行速度要慢得多，如果他們主宰的活動（如幼稚減少或幼稚直方圖化。）一般來說，加快原子操作的方法是不使用它們，或者將它們有節制地以一種不主導計劃活動的方式。例如，天真的縮減會使用一個原子來將每個元素添加到全局和。智能並行壓縮將不會使用線程塊中完成的工作原子。在線程塊減少結束時，可能會使用單個原子將線程塊部分和更新爲全局和。這意味着我可以快速並行地減少任意數量的元素，大概32個或更少的元素。這節約使用原子能的基本上不會在整個程序執行引人注目，但它使平行降低在單一內核調用來完成，而不是2

共享內存：難道他們跑parralel或者他們會排隊？

他們將進行排隊。其原因是可以處理共享內存上的原子操作的功能單元數量有限，不足以在單個週期內處理來自變形的所有請求。

我避免試圖回答與原子操作的吞吐量有關的問題，因爲這些數據在文檔AFAIK中沒有很好地指定。這可能是因爲如果你發出足夠多的同時或幾乎同時的原子操作，由於原子功能單元填滿的隊列已滿，一些warp將停滯在原子指令上。我不知道這是真的，我不能回答有關它的問題。