CUDA通信机制

CUDA体系架构支持怎么样的通信方法呢？往下看喽@_@

1.       __syncthreads()

Block内的线程同步。Block内所有线程都执行到这一位置(BAR指令)，先到的要等后来的，到齐了后再继续后面的任务。执行结果对block内所有线程可见

2.       Memory Fence

l  __threadfence()

Grid内的线程同步。保证该语句前的，grid中所有线程发出的访存指令（对global memory / shared memory）都已结束。执行结果对grid内所有线程可见

l  __threadfence_block()

Block内的线程同步。保证该语句前的，block中所有线程发出的访存指令（对global memory / shared memory）都已结束。执行结果对block内所有线程可见

3.       Mapped pinned memory

CUDA2.2中引入mapped memory，允许多个GPU设备从内核程序中直接访问同一块pinned memory。CPU/GPUs共用一块内存，所以特别要注意做好同步。

4.       cudaThreadSynchronize()

GPU与CPU同步。Kernel启动后控制权异步返回，该函数用以确定所有设备端线程均已运行结束。

5.       volatile

__global__ void myKernel(int* result, int* myArray) {      int tid = threadIdx.x;      int ref1 = myArray[tid] * 1;      myArray[tid + 1] = 2;      int ref2 = myArray[tid] * 1;      result[tid] = ref1 * ref2; }

上面这段程序，myArray[]初值为1。这段程序跑出来的结果result[]全是1，说明计算ref2时myArray[tid]并如有如约变成2（由tid-1修改），程序员是想用改变后的值2来着。为什么会出现这种情况呢？

是这样，在计算ref1时从显存取出myArray[tid]，然后在计算ref2时compiler不会再去显存取一次而是去复用刚刚取到的结果，myArray[tid+1]=2这句操作呢是要写回显存，所以喽^o^ 可以通过加上volatile关键字来解决。

__global__ void myKernel(int* result, volatile int* myArray)

volatile限定符将变量声明为敏感变量，compiler认为其它线程可能随时会修改变量的值，因此每次对该变量的引用都会被编译成一次真实的访存指令。注意，在上面的代码中即使将myArray声明为valatile仍不能确保ref的值是2，这是因为线程tid可能在myArray[tid]被赋为2之前就已经进行了读操作，怎么办呢？做一个同步就好喽。

6.       atomic function

当多个线程同时访问global memory / shared memory同一位置时，原子函数来确保对这个32-/64-bit字的read-modify-write原子操作。保证每个线程能够实现对共享可写数据的互斥操作：在一个操作完成前，其它任何线程都无法访问此地址。

7.       vote

CUDA2.0引入的新特性，1.2计算能力的硬件支持。Vote指令作用域是一个warp。

int __all(int predicate); //warp中所有线程的判断表达式结果为真，则返回1，否则0 int __any(int predicate); //warp中只要一个thread的判断表达式结果为真，则返回1，否则0