http://www.cnblogs.com/dwdxdy/p/3215136.html

__syncthreads()是cuda的內(nèi)建函數(shù)，用于塊內(nèi)線程通信.

__syncthreads() is you garden variety thread barrier. Any thread reaching the barrier waits until all of the other threads in that block also reach it. It is

designed for avoiding race conditions when loading shared memory, and the compiler will not move memory reads/writes around a __syncthreads().

其中，最重要的理解是那些可以到達(dá)__syncthreads()的線程需要其他可以到達(dá)該點的線程，而不是等待塊內(nèi)所有其他線程。

一般使用__syncthreads()程序結(jié)構(gòu)如下：

1 __share__ val[]; 2 ... 3 if(index < n) 4 { 5 if(tid condition) 6 { 7 do something with val; 8 } 9 __syncthreads(); 10 do something with val; 11 __syncthreads(); 12 }

這種結(jié)構(gòu)塊內(nèi)所有線程都會到達(dá)__syncthreads()，塊內(nèi)線程同步.

1 __share__ val[]; 2 ... 3 if(index < n) 4 { 5 if(tid condition) 6 { 7 do something with val; 8 __syncthreads(); 9 } 10 else 11 { 12 do something with val; 13 __syncthreads(); 14 } 15 }

這種結(jié)構(gòu)將塊內(nèi)線程分成兩部分，每一部分對共享存儲器進(jìn)行些操作，并在各自部分里同步.這種結(jié)構(gòu)空易出現(xiàn)的問題是若兩部分都要對某一地址的共享存儲器進(jìn)行寫操作，將可能出

現(xiàn)最后寫的結(jié)果不一致錯誤.要讓錯誤不發(fā)生需要使用原子操作.

1 __share__ val[]; 2 .... 3 if(index < n) 4 { 5 if(tid condition) 6 { 7 do something with val; 8 __syncthreads(); 9 } 10 do something with val; 11 }

這種結(jié)構(gòu)，塊內(nèi)只有部分線程對共享存儲器做處理，并且部分線程是同步.那些不滿足if條件的線程，會直接執(zhí)行后面的語句.若后面的語句里面和if里面的語句都對共享存儲器的同一

地址進(jìn)行寫操作時將會產(chǎn)生wait forever。若沒有這種情況出現(xiàn)，程序則可以正常執(zhí)行完.

在使用if condition 和__syncthreads()，最好使用第一結(jié)構(gòu)，容易理解，不容易出錯

總結(jié)

以上是生活随笔為你收集整理的__syncthreads()的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網(wǎng)站內(nèi)容還不錯，歡迎將生活随笔推薦給好友。