【IT168 文档】CUDA体系架构支持怎么样的通信方法呢?往下看喽@_@

　　1. __syncthreads()

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

　　2. Memory Fence

• __threadfence()

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

• __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　　上面这段程序，myArray[]初值为1。这段程序跑出来的结果result[]全是1，说明计算ref2时myArray[tid]并如有如约变成2(由tid-1修改)，程序员是想用改变后的值2来着。为什么会出现这种情况呢?

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

 　　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。