对__threadfence的一点理解

一直没搞清楚,cuda 2.2版增加的__threadfence到底有何作用,直到今天看到sdk 3.0手册
中的下面例子才恍然大悟.(中文为我的理解,嘿嘿)

 

一个求和的例子:
__device__ unsigned int count = 0;// 统计有几个block结束的变量
__shared__ bool isLastBlockDone;  // 第一轮(多block)计算是否结束
__global__ void sum(const float* array, unsigned int N, float* result)
{
   // Each block sums a subset of the input array
   float partialSum = calculatePartialSum(array, N);
      // 第一轮计算,多个block参加
      // calculatePartialSum中应该有__syncthreads同步
   if (threadIdx.x == 0) // 参加的第一轮计算的每个block计算完毕后最后用线程0保存
                               // 本block计算结果并参与竞争第二轮计算的资格
   {
      // Thread 0 of each block stores the partial sum to global memory
      result[blockIdx.x] = partialSum;
      // Thread 0 makes sure its result is visible to
      // all other threads
      __threadfence();
      // 保证result[blockIdx.x] = partialSum已经执行完毕了!
      // 否则,cuda是写完不管的,可能会有如下情况出现:
      // block1计算结束写保证result....block2计算结束写保证result...block2获得第二轮
      //  计算资格......block2写result完成...block2读result(第二轮计算).... block1写result
      //  完成...block1结束....
      // 而加了__threadfence以后,在block1执行下面的atomicInc以前,block1的result写
      // 已经完成了!因此，无论是哪个block最后获得第二轮计算权，前面的result肯定
      // 都已经完成。这个就是__threadfence的意义，否则，我们只有用kernel间同步了
     
      // Thread 0 of each block signals that it is done
      unsigned int value = atomicInc(&count, gridDim.x);
      // Thread 0 of each block determines if its block is
      // the last block to be done
      isLastBlockDone = (value == (gridDim.x - 1));
   }  // Synchronize to make sure that each thread reads
   // the correct value of isLastBlockDone
   __syncthreads();
   // 因为不确定是哪个block会获得第二轮计算权，因此，block中非0的其它线程要
   // 在这里同步等待线程0的竞争结果。
   if(isLastBlockDone)
   {  // isLastBlockDone是smem变量，每个block是不同的!而只有对一个block，它的
      // 值会为真
     
      // The last block sums the partial sums
      // stored in result[0 .. gridDim.x-1]
      float totalSum = calculateTotalSum(result);
      if (threadIdx.x == 0)
      {
          // Thread 0 of last block stores total sum
          // to global memory and resets count so that
          // next kernel call works properly
          result[0] = totalSum;
          // 当kernel结束时，gmem写也自然结束了。因此，这里不要__threadfence了。呵呵。
          count = 0;
      }
   }
   // 其它的block就直接结束了
}

 

个人感觉，__threadfence一般用于block间有竞争且竞争成功后要用到其它block的前面的
全局写的结果的场合比较合适(当然，也可以用拆kernel的方法)。而block内还是直接用
__syncthreads更直观些。。。。

 

一句话总结，有些鸡肋的味道。