之前在寫程式的時候，經常用弄混同步函式，現做出總結。
_syncthreads（）：執行緒塊內線程同步；保證執行緒會腫的所有執行緒都執行到同一位置； 當整個執行緒塊走向同一分支時才可以使用_syncthreads()，否則造成錯誤；乙個warp內的執行緒不需要同步；即當執行的執行緒數小於warpsize時，不需要同步函式，呼叫一次至少需要四個時鐘週期，一般需要更多時鐘週期，應盡量避免使用。
每個sm包含8個cuda核心，並且在任何乙個時刻執行32個執行緒的單個warp ， 因此需要4個時鐘週期來為整個warp發布單個指令。你可以假設任何給定warp中的執行緒在鎖步（
lockstep
）中執行，。lockstep技術可以保持多個cpu、記憶體精確的同步，在正確的相同
時鐘週期
內執行相同的指令。但要跨warp進行同步，您需要使用 __ syncthreads（）。
這裡主要區別三個同步函式：cudastreamsynchronize、cudadevicesynchronize 和 cudathreadsynchronize。在文件中，這三個函式叫做barriers，只有滿足一定的條件後，才能通過barriers向後執行。三者的區別如下：
cudadevicesynchronize()
：該方法將停止cpu端線程的執行，直到gpu端完成之前cuda的任務，包括kernel函式、資料拷貝等。
cudathreadsynchronize()
：該方法的作用和cudadevicesynchronize()基本相同，但它不是乙個被推薦的方法，也許在後期版本的cuda中會被刪除。
cudastreamsynchronize()
：這個方法接受乙個stream id，它將阻止cpu執行直到gpu端完成相應stream id的所有cuda任務，但其它stream中的cuda任務可能執行完也可能沒有執行完。
在cuda裡面，不同執行緒間的資料讀寫會彼此影響，這種影響的作用效果根據不同的執行緒組織單位和不同的讀寫物件是不同。
在不考慮2.x的優化的情況下，
(1)在同乙個warp內的執行緒讀寫shared/global,
讀寫global和shared是立刻對本warp內的其他執行緒立刻可見的。
(2)在同乙個block內的不同warp內線程讀寫shared/global,
這種讀寫必須使用__syncthreads(), 或者__threadfence()來實現不同的讀寫可見效果。
(3)在同乙個grid內的不同block內的執行緒讀寫shared/gloabl,
這種讀寫必須使用__threadfence*()來實現一定的讀寫可見效果。
(4)任何執行緒組織單位內的原子操作總是可見的。
所以:
樓主寫到:if(threadidx.x == 0) //樓主假設了block形狀是1維的
__syncthreads();
if(islastblockdone)
}
樓主還寫到:這種說法的本意是對的，但是說的不對。這本書的意思我想應該是，書上說：如果在儲存部分和與計數器遞增只講不進行fence操作，則計數器的值可能在儲存子串行之和之前就已經遞增了（這個情況怎麼會發生，計數器加1不是在result賦值之後麼？）。此時最後乙個block讀到的子串行和可能還沒更新，造成計算結果錯誤。
如果在儲存部分和與計數器遞增之間不使用__threadfence()，則計數器的值可能， 在儲存子串行之和對所有的grid內線程都可見之前就已經遞增了。您所說的"這個情況怎麼會發生，計數器加1不是在result賦值之後麼？", 前面的寫入中間結果的操作的確是在atomicinc()之前，但是他們的對別的執行緒單位的讀寫效果可見性，生效時間不是同的。
也就是說，因為您的**是這個grid內的最後乙個block在進行中間結果加和。
而同時，因為atomicinc的讀寫效果是立即的，如果您這裡沒有使用__threadfence(), 那麼當進入了最後乙個block，開始求和的時候，由於前面的儲存中間結果的其他block的寫入效果，對於本最後乙個block中的執行緒來說，沒有立刻生效。
也就是說，最後block中的執行緒有一定可能看不到這種寫入的值。所以您必須在讓最後乙個block之前，用__threadfence()來確保所有的block中的第0個執行緒的寫入效果全grid都可見。而您的**，這種地方只能寫在atomicinc之前了
。總之，讀寫是你看到的，但是生效是隱藏在後面的。
可能會覺得有程式中有乙個__syncthreads();這樣的話，似乎那個__threadfence（）顯得多餘。但是：
根據手冊的說法。__syncthreads()可以讓同乙個block的執行緒們在這個點同步，同時的附加效果是，本block內的所有執行緒的讀寫對彼此立刻有效。
看到您後面是不同的block得到的值，由最後乙個block來求總和。所以我覺得只有乙個__syncthreads()是不夠的。
__syncthreads()只保證了對塊內的可見，並保證沒有對其他塊的可見，threadfence保證了這點。
摘自：悠閒的小貓的解答。
CUDA之Thread Wrap執行詳解
從硬體角度分析，支援cuda的nvidia 顯示卡，都是由多個multiprocessors 組成。每個 multiprocessor 裡包含了8個stream processors，其組成是四個四個一組，也就是兩組4d的處理器。每個 multiprocessor 還具有 很多個 比如8192個 暫...
CUDA之Thread Wrap執行詳解
2017 03 23 13 53 1052人閱讀收藏 舉報cuda 32 從硬體角度分析，支援cuda的nvidia 顯示卡，都是由多個multiprocessors 組成。每個 multiprocessor 裡包含了8個stream processors，其組成是四個四個一組，也就是兩組4d的處理...
cuda 矩陣乘法函式之cublasSgemm
可以考慮使用,例如cublassgeam 矩陣加法 進行一次1.0 at 0.0 b的引數設定,利用內建的轉置功能 注意這裡的1和0 來進行將a轉換成at.在使用cuda的cublas庫中矩陣乘法函式cublassgemm時，注意到cuda其中的二維矩陣的儲存是 按列儲存 一天都處於蒙蔽狀態，查了很...