关于CUDA中_threadfence的理解

     __threadfence函数是memory fence函数，用来保证线程间数据通信的可靠性。与同步函数不同，memory fence不能保证所有线程运行到同一位置，只保证执行memory fence函数的线程生产的数据能够安全地被其他线程消费。

（1）__threadfence：一个线程调用__threadfence后，该线程在该语句前对全局存储器或共享存储器的访问已经全部完成，执行结果对grid中的所有线程可见。

（2）__threadfence_block：一个线程调用__threadfence_block后，该线程在该语句前对全局存储器或者共享存储器的访问已经全部完成，执行结果对block中的所有线程可见。

     上面是官方解释，但是看完之后还是不明其所以然，尤其是手册中给的元素求和代码。为了明白其真正函数，自己把代码实现了一遍，然后通过运行明白了__threadfence的含义与作用。

__device__ int count=0;__global__ static void sum(int* data_gpu,int* block_gpu,int *sum_gpu,int length){    extern __shared__ int blocksum[];    __shared__ int islast;    int offset;    const int tid=threadIdx.x;    const int bid=blockIdx.x;    blocksum[tid]=0;    for(int i=bid*THREAD_NUM+tid;i<length;i+=BLOCK_NUM*THREAD_NUM)    {        blocksum[tid]+=data_gpu[i];    }     __syncthreads();    offset=THREAD_NUM/2;    while(offset>0)    {        if(tid<offset)        {            blocksum[tid]+=blocksum[tid+offset];        }        offset>>=1;        __syncthreads();    }    if(tid==0)    {        block_gpu[bid]=blocksum[0];       __threadfence();        int value=atomicAdd(&count,1);        islast=(value==gridDim.x-1);    }    __syncthreads();    if(islast)    {        if(tid==0)        {            int s=0;            for(int i=0;i<BLOCK_NUM;i++)            {                s+=block_gpu[i];            }            *sum_gpu=s;        }    }}

      上述CUDA代码实现了block之间对元素求和，关键的地方在32行和34行（手册貌似有错误，用的是gridDim.x，我改成了1）。起初自己认为__threadfence或者原子操作单独都可以完成运算，所以我通过分别去除__threadfence和后面的原子操作来验证结果的正确性，结果发现：

• 单独的__threadfence不能给出正确结果；
• 只用原子操作可以给出正确结果。

       一开始对结果很奇怪，然后从网上搜各种解释，得到的结论是：threadfence不是保证所有线程都完成同一操作，而只保证正在进行fence的线程本身的操作能够对所有线程安全可见fence不要求线程运行到同一指令，而barrier有要求。上述结论指出__threadfence函数不是同步函数，如果单纯地让block 0去计算最终的结果，这时可能会存在还有其他block尚未执行，这时得到的结果必然是错误的。

     虽然只用原子操作可以给出正确结果，但是也不能保证在其他情况下也是正确的（GPU编程需要特别注意当前条件下正确的程序换个条件不一定正确，反映了GPU编程的复杂性）。这里正确的原因可能是因为访问的全局内存只有一个空间，原子操作也是访问全局内存中的变量，这两个访问时间属于一个量级导致，如果一开始的访问全局内存不是一个空间，而是一个比较长的数组，则此时就可能会出错。为什么会出错，这和CUDA对全局内存的读写有关：线程在读取全局存储器的时候会被阻塞，然后warp scheduler会接着调度其他warp；但是当线程在写入全局内存时，虽说该写入操作尚未完成，但是线程会接着执行下面的指令，而不是等待写入完成。在这种情况下，如果访问不同block中的数据，不加__threadfence确实会存在出错的可能。

      现在重新考虑一种情况，在原子操作之前的全局内存方法写入的是一个长数组，然后我们去掉__threadfence，只用原子操作来保证正确性。此时最后对原子变量进行操作的block完成之后开始对全局内存进行读操作。由于block调度的不确定性，这时可能会存在其他block中的线程尚未完成全局内存的写入，此时访问其他block要写入的全局内存就会出错。虽然本block内的全局内存也有可能尚未完成，但是同一warp内会有写后读的限制，同时同一block可以通过__syncthreads来同步，只有不同block之间的全局内存访问不能保证，所以__threadfence是必须的。

    关于数据在不同线程间的可见性具有下面的结论：

1. 在同一个warp内的线程读写shared/global,读写global和shared是立刻对本warp内的其他线程立刻可见的。
2. 在同一个block内的不同warp内线程读写shared/global,这种读写必须使用__syncthreads(), 或者__threadfence_block()来实现不同的读写可见效果。
3. 在同一个grid内的不同block内的线程读写shared/gloabl, 这种读写必须使用__threadfence()来实现一定的读写可见效果。
4. 任何线程组织单位内的原子操作总是可见的。

    总结下，block之间的内存访问模型可以简化为以下三步：

1. store your data
2. __threadfence()
3. atomically mark a flag

遵循上面的步骤，既可以保证正确性，又可以获得一个较高的速度。

   可以参考的资料有：

1. 对__threadfence的一点理解
2. 关于_threadfence的问题归约求和程序，每次运行结果都不相同 
3. 归约求和程序，每次运行结果都不相同 
4. CUDA __threadfence()
   
5. Inter-Block GPU Communication via Fast Barrier Synchronization
   
6. To GPU Synchronize or Not GPU Synchronize?