__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()，最好使用第一结构，容易理解，不容易出错～