【CUDA學習】__syncthreads的理解

__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()程序結構如下：

__share__ val[];
...
if(index < n)
{    
   if(tid condition)    
    {         
        do something with val;    
    }    
    __syncthreads();    
    do something with val;
    __syncthreads();
}    

這種結構塊內所有線程都會到達__syncthreads()，塊內線程同步.

__share__ val[];
...
if(index < n)
{     
    if(tid condition)     
    {          
        do something with val;
        __syncthreads();      
    }     
    else      
    {         
        do something with val;
        __syncthreads();     
    }
}            

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

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

__share__ val[];
....
if(index < n)
{      
    if(tid condition)      
    {          
        do something  with val;          
        __syncthreads();       
    }      
    do something with val;
}

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

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

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