cuda基础---cuda通信机制

 1 __device__ unsigned int count = 0;
 2 __shared__ bool isLastBlockDone;
 3 
 4 //每个ｂｌｏｃｋ首先计算出数组的一个子序列的和，再将结果存储到全局内存
 5 //所有ｂｌｏｃｋ完成后，由最后一个块完成求和操作
 6 
 7 __global__　void sum(const float* array,unsigned int N,float* result)
 8 {
 9 //每个ｂｌｏｃｋ对输入数组的一个ｓｕｂｓｅｔ求和
10 float partialSum = calcultePartialSum(array,N);
11 if(threadIdx.x == 0)
12 {
13 //０号线程存储局部和ｐａｒｔｉａｌＳｕｍ并负责写回全局存储器
14 result[blockIdx.x] = partialSum;
15 //线程０要保证它的结果对所有其他的线程可见
16 __threadfence();
17 //每个ｂｌｏｃｋ的０号线程负责计算结束后做一次标记，ｃｏｕｎｔ＋＋
18 unsigned int value = atomicInc(&count,gridDim.x);
19 //每个ｂｌｏｃｋ的０号线程来判断该块是否是最后一个计算结束的快
20 isLastBlckDone = (value == gridDim.x-1);
21 }
22 //做一次同步，保证每个线程读到正确的ｉｓＬａｓｔＢｌｏｃｋＤｏｎｅ值
23 __syncthreads();
24 
25 if(isLastBlockDone)
26 {
27 //最后的那个ｂｌｏｃｋ负责局部和的球和
28 float totalSum = calculateTotalSum(result);
29 if(threadIdx.x == 0)
30 {／／ｌａｓｔ　ｂｌｏｃｋ的线程０存储急速那和并将其协会全局存储器，ｃｏｕｎｔ　置０来保证下一个ｋｅｒｎｅｌ的正常计数
31 result[0] = totalSum;
32 count = 0;
33 }
34 }
35 }

 1 __global__ void testAtom(int* g_odata)
 2 {
 3 const unsigned int tid = blockDim.x*blockIdx.x+threadIdx.x;
 4 
 5 //算术原子操作
 6 atomicAdd(&g_odata[0],10);
 7 atomicMax(&g_odata[0],tid);
 8 atomicInc((unsigned int*)&g_odata[0],17);
 9 atomicCAS(&g_odata[0],tid-1,tid);
10 
11 //位原子操作
12 atomicAnd(&g_odata[8],2*tid+7);
13 atomicOr(&g_odata[8],1<<tid);
14 atomicXor(&g_odata[8],tid);        
15 }