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；用于判断线程活跃性；

测量程序运行时间：

CUDA内核程序的运行时间：可以在设备端测量，也可以在主机端测量；

CUDA API的运行时间：只能在主机端测量；使用CUDA runtime API时，会在第一次调用runtime API函数时启动CUDA环境，计时的时候应该避免将这一部分计入，因此在正式测试之前应当首先及你选哪个一侧包含数据输入输出地就爱上你，使得GPU从平时的节能模式进入工作状态，使得测试结果更加可靠；

 

设备端测量时间：

调用clock（）函数：返回的是GPU的时钟周期，需要除以GPU的运行频率才能得到以秒为单位的时间；

使用CUDA API事件管理功能；

 

主机端测量时间：

使用c标准库中的clock_t()函数测试，由于其精度很低，因此应该运行多次然后求平均运行时间；注意异步函数（比如内核函数和带有asyn后缀的存储器拷贝函数），在GPU上执行完成之前，CPU线程已经得到了它的返回值；从主机测量一系列CUDA调用需要的时间的时候，要首先调用cudaThreadSynchronize()函数等，使得GPU线程执行完毕后，进入CPU线程，从而得到正确的执行效果；在一串流中的第一个流（ID为0的流）的行为总是同步的，因此使用这些函数对0号流进行测时，得到的记过是可靠的。

 

转自：http://blog.csdn.net/ouczoe/archive/2010/01/03/5125694.aspx