__syncthreads()

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

__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 }