CUDA之同步函数详解

之前在写程序的时候，经常用弄混同步函数，现做出总结。

_syncthreads（）：线程块内线程同步；保证线程会肿的所有线程都执行到同一位置； 当整个线程块走向同一分支时才可以使用_syncthreads()，否则造成错误；一个warp内的线程不需要同步；即当执行的线程数小于warpsize时，不需要同步函数，调用一次至少需要四个时钟周期，一般需要更多时钟周期，应尽量避免使用。每个SM包含8个CUDA内核，并且在任何一个时刻执行32个线程的单个warp ， 因此需要4个时钟周期来为整个warp发布单个指令。你可以假设任何给定warp中的线程在锁步（LOCKSTEP）中执行，。LOCKSTEP技术可以保持多个CPU、内存精确的同步，在正确的相同时钟周期内执行相同的指令。但要跨warp进行同步，您需要使用 __ syncthreads（）。

        这里主要区别三个同步函数：cudaStreamSynchronize、CudaDeviceSynchronize 和 cudaThreadSynchronize。在文档中，这三个函数叫做barriers，只有满足一定的条件后，才能通过barriers向后执行。三者的区别如下：
         cudaDeviceSynchronize() ：该方法将停止CPU端线程的执行，直到GPU端完成之前CUDA的任务，包括kernel函数、数据拷贝等。
         cudaThreadSynchronize() ：该方法的作用和cudaDeviceSynchronize()基本相同，但它不是一个被推荐的方法，也许在后期版本的CUDA中会被删除。

        cudaStreamSynchronize()：这个方法接受一个stream ID，它将阻止CPU执行直到GPU端完成相应stream ID的所有CUDA任务，但其它stream中的CUDA任务可能执行完也可能没有执行完。

在CUDA里面，不同线程间的数据读写会彼此影响，这种影响的作用效果根据不同的线程组织单位和不同的读写对象是不同。  

在不考虑2.x的优化的情况下，  

(1)在同一个warp内的线程读写shared/global,  
  读写global和shared是立刻对本warp内的其他线程立刻可见的。  

(2)在同一个block内的不同warp内线程读写shared/global,  
    这种读写必须使用__syncthreads(), 或者__threadfence()来实现不同的读写可见效果。  

(3)在同一个grid内的不同block内的线程读写shared/gloabl,  
    这种读写必须使用__threadfence*()来实现一定的读写可见效果。  

(4)任何线程组织单位内的原子操作总是可见的。  
所以:  

楼主写到: 
if(threadIdx.x == 0) //楼主假设了block形状是1维的 
{        result[blockIdx.x] = partialSum; 
            __threadfence(); 
            unsigned int value = atomicInc(&count ,gridDim.x); 
             isLastBlockDone = (value ==griddim.x-1) 
} 
  __syncthreads(); 
if(isLastBlockDone) 
{ 
      floa totalSum = calclateTotalSum(rusult); 
      if(threadIdx.x ==0) 
      { 
            result[0] = totalSum; 
            count = 0; 
      } 
}

楼主还写到: 
书上说：如果在存储部分和与计数器递增只讲不进行fence操作，则计数器的值可能在存储子序列之和之前就已经递增了（这个情况怎么会发生，计数器加1不是在result赋值之后么？）。此时最后一个block读到的子序列和可能还没更新，造成计算结果错误。

这种说法的本意是对的，但是说的不对。这本书的意思我想应该是，   如果在存储部分和与计数器递增之间不使用__threadfence()，则计数器的值可能， 在存储子序列之和对所有的grid内线程都可见之前就已经递增了。您所说的"这个情况怎么会发生，计数器加1不是在result赋值之后么？",  前面的写入中间结果的操作的确是在atomicInc()之前，但是他们的对别的线程单位的读写效果可见性，生效时间不是同的。

也就是说，因为您的代码是这个grid内的最后一个block在进行中间结果加和。
而同时，因为atomicInc的读写效果是立即的，如果您这里没有使用__threadfence(), 那么当进入了最后一个block，开始求和的时候，由于前面的存储中间结果的其他block的写入效果，对于本最后一个block中的线程来说，没有立刻生效。

也就是说，最后block中的线程有一定可能看不到这种写入的值。所以您必须在让最后一个block之前，用__threadfence()来确保所有的block中的第0个线程的写入效果全grid都可见。而您的代码，这种地方只能写在atomicInc之前了   。  

总之，读写是你看到的，但是生效是隐藏在后面的。

可能会觉得有程序中有一个__syncthreads();这样的话，似乎那个__threadfence（）显得多余。但是：

根据手册的说法。__syncthreads()可以让同一个block的线程们在这个点同步，同时的附加效果是，本block内的所有线程的读写对彼此立刻有效。

看到您后面是不同的block得到的值，由最后一个block来求总和。所以我觉得只有一个__syncthreads()是不够的。

__syncthreads()只保证了对块内的可见，并保证没有对其他块的可见，threadfence保证了这点。

摘自：http://hpcbbs.it168.com/thread-1818-1-1.html

悠闲的小猫的解答。

http://www.voidcn.com/blog/hhko12322/article/p-1947268.html