CUDA通信机制之同步函数（二）

[_梦魇花葬]

上次我们讨论了关于blokc内线程同步，这次我们来探讨一下block间线程同步和kernel间同步。
同步函数：memory fence 函数
memory fence 函数是用来保证线程间数据通信的可靠性的，但与__syncthreads()函数不同的是，memory fence 函数并不要求所有线程都运行到同一位置，它只保证执行memory fence 函数的线程生产的数据能够安全地被其他线程消费。
(1) __threadfence()：一个线程调用__threadfence()后，该线程在该语句前对全局存储器或共享存储器的访问已经全部完成，执行的结果对grid中的所有线程可见。
(2) __threadfence_block()：一个线程调用__threadfence_block()后，该线程在该语句前对全局存储器或共享存储器的访问已经全部完成，执行的结果对block中的所有线程可见。
memory fence函数使其他线程能够安全地消费当前线程生产的数据，多个线程间可以正确地操作共享数据，实现grid/block内的线程间通信。比如下面的例子：
__device__ unsigned int count = 0;
__shared__ bool isLastBlockDone;
__global__ void sum(const float* array, unsigned int N, float* result){
//每个block对输入数组的一个subset求和
float partialSum = calculatePartialSum(array, N);
if(threadIdx.x == 0){
//0号线程存储局部和partialSum并负责写回全局存储器
result[blockIdx.x] = partialSum;
// 线程0要保证它的结果对所有其他的线程可见
__threadfence()；
// 每个block的0号线程负责计算结束后做一次标记，count++
unsigned int value = atomicInc(& count, gridDim.x);
// 每个block的0号线程来判断该块是否是最后一个计算结束的块
isLastBlockDone = (vlaue == (gridDim.x - 1));
}
//做一次同步，保证每个线程读到正确的isLastBlockDone 值
__syncthreads();

if(isLastBlockDone){
// 最后的那个block负责局部和（存储在result[0 .. gridDim.x -1]）的求和
float totalSum = calculateTotalSum(result);
if(threadIdx.x == 0){
// last block的线程0存储计算总和并将其写回全局存储器，count置0以保证下一个kernel的正常计数
result[0] = totalSum;
count = 0;
}
}
}
上面的例子实现的是对N个元素求和：每个block首先计算出数组的一个子序列的和，然后将结果存储在全局存储器中，当所有block完成这一步以后，再由最后一个完成求和操作的block从全局存储器中读入所有子序列之和，并计算最终结果。如果在存储部分和与计数器递增之间不进行fence操作，那么计数器中的值可能在存储子序列和之前就已经递增了，而此时最后一个block读到得子序列和可能还没有更新，造成计算结果错误。
kernel 间通信：可以同过global memory实现。GPU间的通信需要通过主机端内存进行，代价高昂；mapped memory功能允许多个设备从内核程序中直接访问同一块pinned memory。

[伊顺鸣]

看看的啊，，，，，

[_梦魇花葬]

恩，2楼有什么问题可以提出来的哟！~~~一起成长！