CUDA学习（六）

CUDA线程执行模型为两层，一层Grid,一层Block，block之间时不可以通信的，但是block里面的线程可以相互通信，逻辑层和物理层的实现，以及代码端的书写都需要好好学习

• __syncthreads()

__syncthreads()实现了线程块内的线程同步，它保证了线程块中的所有线程都执行到同一位置。这都是线程操作所必须的内容。

只有当整个线程块走向相同分支时，才能在条件语句里面使用__syncthreads(),否则可能引起错误。另外，一个warp内的线程不用同步，也就是说，如果需要同步的线程处于同一warp中，则不需要调用__syncthreads()。

• memory fence函数

memory fence函数用来保证线程生产的数据能够安全地被其他线程消费。

__threadfence():该线程在该语句前对全局存储器或共享存储器的访问已经全部完成，执行的结果对grid中的所有线程可见。

__threadfence_block():该线程在该语句前对全局存储器或共享存储器的访问已经全部完成，执行的结果对block中的所有线程可见。

• kernel间通信

kernel之间的数据传递，可以通过global memory实现

而GPU之间的通信主要依靠mapped memory实现

• GPU与CPU之间的线程同步

在CUDA主机端代码中使用cudaThreadSynchronize()，可以实现GPU与CPU线程的同步。kernel启动后控制权将异步返回，利用该函数可以确定所有设备端线程均已运行结束。

• Volatile关键字

Volatile关键字与C++的该关键字具有相同的作用，都时保证自己的代码不被优化，每一行相关的代码都被编译为指令

• Atom操作

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