主要涉及三个函数 和原子同步指令
1 __syncthreads(); 使得同一个block之间线程间同步，达到相同的执行点后再往后执行，同时使得修改的全局以及共享内存对block内的线程可见
2 __threadfence(); 该线程在该语句前对全局存储器或共享存储器的访问已经全部完成，执行结果对grid中的所有线程可见。
3 __thradfence_block(); 该线程在该语句前对全局存储器或者共享存储器的访问已经全部完成，执行结果对block中的所有线程可见。
注意CUDA中一个GRID包含多个BLOCK。每个BLOCK包含多个THREAD。另外还有一个概念叫做wrap。WRAP是线程束的意思，也就是GPU实际执行运算的时候是以wrap单位的。比如wrap=32或者wrap=16.假设你的wrap是16，每次运行一定是16个线程一起运行。即使你使用了一个block里面的1个thread.GPU也会凑足16个thread，只是这些thread处于不活跃状态，此时就浪费了15个thread线程的资源。
http://www.myexception.cn/cuda/1931284.html
wrap了解清楚以后再看内存同步解释
1 同一个wrap内,对于global/shared 的内存操作是 该wrap内的所有线程即时可见的。
例如某个线程修改了一个内存，那么wrap内的其他线程访问的这个内存是修改后的。
2同一个block内。对于global/shared的内存操作。block内的线程不一定是即时可见的。
例如一个block内有 32个线程。32个线程属于两个wrap wrapA wrapB。wrapA修改了某个内存，wrapB里面访问不一定是修改后的（wrapA访问的一定是修改后的，见第一条）。因此需要使用
__threadfence_block() _syncthread() __threadfence()来实现让修改数据对其他线程可见。
3同一个grid内，对于global/shard的内存操作。其他线程不一定是可见的
必须使用__threadfence()来实现让修改数据对其他线程可见。其他两个函数只对于block有效。
4原子操作。atomic的操作修改对于所有的线程都是即时可见的。
比如使用多个block对一个向量进行求和运算。向量长为30。分为三个block
a[0]-------a[29]
首先 每个block分别计算10个数相加。block0 计算 a[0]-a[9] block1 计算 a[10]-a[19] block2 计算 a[20]-a[29]。计算的结果放到a[blockIdx.x]里面。
由最后一个线程计算 a[0]+a[1]+a[2]得到最后结果。如何知道是最后一个block呢？我们用一个count来计数。
每个block的线程0计算完毕的时候。使用原子加 atomicadd使得count+1.如果count等于block的大小。那么说明是最后一个block了。可以由他来执行最后的a[0]+a[1]+a[2]了
此时就需要用到_threadfence()确保。a[1]和a[2]的结果对于block0的最后计算线程是可见的。
if(threadIdx.x == 0)
{
执行fence保证a[1]和a[2]的结果对于后续的 float totalSum = calclateTotalSum(rusult);是可见的。
__threadfence();
unsigned int value = atomicInc(&count ,gridDim.x);
isLastBlockDone = (value ==griddim.x-1)
}
__syncthreads();
if(isLastBlockDone)
{
float totalSum = calclateTotalSum(rusult);
if(threadIdx.x ==0)
{
result[0] = totalSum;
count = 0;
}
}
还有一个函数就走__threadfence_system()。这个就是保证所有的内存读写对于所有线程已完全可见了。包括主机线程等等
关于内存可以参见 http://www.mamicode.com/info-detail-226119.html
