一直没搞清楚,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更直观些。。。。
一句话总结,有些鸡肋的味道。