【CUDA學習】__syncthreads的理解

一點心青發表於2013-07-25
thread

__syncthreads()是cuda的內建函式，用於塊內執行緒通訊.

__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().

其中，最重要的理解是那些可以到達__syncthreads()的執行緒需要其他可以到達該點的執行緒，而不是等待塊內所有其他執行緒。

一般使用__syncthreads()程式結構如下：

 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 }

這種結構塊內所有執行緒都會到達__syncthreads()，塊內執行緒同步.

 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 }

這種結構將塊內執行緒分成兩部分，每一部分對共享儲存器進行些操作，並在各自部分裡同步.這種結構空易出現的問題是若兩部分都要對某一地址的共享儲存器進行寫操作，將可能出

現最後寫的結果不一致錯誤.要讓錯誤不發生需要使用原子操作.

 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 }

這種結構，塊內只有部分執行緒對共享儲存器做處理，並且部分執行緒是同步.那些不滿足if條件的執行緒，會直接執行後面的語句.若後面的語句裡面和if裡面的語句都對共享儲存器的同一

地址進行寫操作時將會產生wait forever。若沒有這種情況出現，程式則可以正常執行完.

在使用if condition 和__syncthreads()，最好使用第一結構，容易理解，不容易出錯～

相關文章