CUDA学习笔记之同步操作指令

同步函数：

_syncthreads（）：线程块内线程同步；

                           保证线程会肿的所有线程都执行到同一位置；

                           当整个线程块走向同一分支时才可以使用_syncthreads()，否则造成错误；

                           一个warp内的线程不需要同步；

                          调用一次至少需要四个时钟周期，一般需要更多时钟周期，应尽量避免使用；

memory fence：保证线程间数据通信的可靠性；

                       不保证所有线程运行到同一位置，之保证执行memory fence函数的线程生产的数据能够安全得被其他线程消费；

                        _threadfence():一个线程调用该函数后，该线程在该语句钱对全局存储器或者共享存储器的访问已经全部完成，执行结果对grid中的所有线程可见；

                       _threadfence_block():一个线程调用该函数后，该线程在该语句钱对全局存储器或共享存储器的访问已经全部完成，执行结果对block中所有线程可见；

                        事实上起到的作用是，及时通知其他线程，全局存储器和共享存储器的结果已经改变了；

 

kernel间的通信：

               可以同过global memory实现。GPU间的通信需要通过主机端内存进行，代价高昂；mapped memory功能允许多个设备从内核程序中直接访问同一块pinned memory。

 

GPU与CPU之间的线程同步：

                主机端代码中使用cudaThreadSynchronize():实现CPU和GPU线程同步；

                kernel启动后控制权将异步返回，利用该函数可以确定所有设备端线程均已运行结束；

                类似的有：cudaStreamSynchronize()和cudaEventSynchronize();

 

Volatile: 敏感变量关键字；编译器认为其他线程可能随时会修改变量的值，因此每次对该变量的引用都会被编译成一次正式的内存读取指令；

 

ATOM操作：互斥变量；对位于全局或共享存储器的一个32为或64为字执行read-modify-write的院子操作；保证每个线程能够实现对共享可写数据的互斥操作；各种硬件对ATOM指令的支持、以及ATOM指令支持的数据类型等不尽相同；

 

VOTE操作： 是cuda2.0的新特性；只有1.2以上版本的硬件才支持；指令作用范围是一个warp；用于判断线程活跃性；