cuda基础---cuda通信机制

1：同步函数

1）__syncthreads()

　　实现线程块中的线程同步----保证线程块中所有线程执行到同一位置;

　　只有当整个线程块都走向相同分支时，才能在条件语句中使用__syncthreads)(;

　　一个warp内的线程不用同步；

　　__syncthreads()调用花费时间周期，应尽量避免/节约使用。

2）memory fence 函数

　　保证数据通信的可靠性，并不要求所有线程都运行到同一位置；

　　多个线程间可以正确的操作共享数据，实现grid/block内的线程间通信。

（1）__threadfence()

　　该函数调用后，该线程在此语句前对全局存储器或共享存储器的访问已经全部完成，且结果对grid内所有线程可见。

（2）__threadfence_block()

　　该函数调用后，该线程在此语句前对全局存储器或共享存储器的访问已经全部完成，且结果对block内所有线程可见。

（3）实例---N个元素求和

 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 }

３）kernel间通信

　　通过global memory实现－－pinned memory

4）ＧＰＵ与ＣＰＵ线程同步

　　Ａ：主机端代码使用cudaThreadSynchronize()----确定所有设备端线程均已运行结束

２：volatile关键字

　　存在于全局变量或共享存储器中的变量通过此关键字声明为敏感变量。每次对该变量的引用都会被编译成一次真实的内存读指令。

３：ATOM操作

　　当多个线程同时访问全局或共享存储器的同一位置时，保证每个线程能够实现对共享可写数据的互斥操作－－在一个操作完成之前，其他任何线程都无法访问此地址。

 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 }

Ａ：图像亮度直方图统计

４：VOTE操作

　　由下面两个函数实现此指令

Ａ：int __all(int predicate);

　　一个ｗａｒｐ中所有的threads的判断表达式结果为真则返回１；否则返回０

Ｂ：int __any(int predicate);

　　一个warp中只要存在ｔhread的判断表达式结果为真，则返回１，否则返回０．