CUDA C编程权威指南 第五章 共享内存和常量内存

1. 共享内存是较小的片上内存,具有较低的延迟(相比全局,低20~30倍),提供更高的带宽(相比全局,10倍)

2. 1. block通信
   2. 用于全局内存数据的缓存

3. __shared__来申请共享变量,如果共享内存大小在编译时是为知的,则需要加载extern关键字extern __shared__ int tile[]

4. 共享内存被分为32个同样大小的内存模型,称为存储体)(bank),有32个bank是因为有32个thread

5. 通过warp共享内存的加载或存储操作,且在每个bank早上只访问不多余一个的内存地址,那么这个操作可以由一个内存事务来完成

6. 共享内存中多个地址请求落在同一个bank中时,会发生内存提冲突,导致请求被重复执行.

单重典型warp发出共享内存请求:
7. 并行访问:多个地址访问多个bank:一个warp访问多个地址落在多个存储体中

8. 串行访问:多个地址访问一个bank:多个地址属于同一个bank,如果warp中的32个thread访问同一个存储体不同地址,就会有32个内存事务
9. 广播访问:单一thread读取一个bank:warp中所有的thread读取同一个bank相同的地址,若一个内存事务被执行,name被访问的地址就会广播到所有请求的线程中,只需要一个广播访问,但是带宽利用率很差,因为海藻油一部分字节被访问

同步

1. 当不同步的多个线程修改同一个共享内存地址时,将导致线程内冲突,cuda提供障碍和内存栅栏来实现同步
2. void __syncthreads()实现障碍,要求block中所有的线程等待直到所有线程到大该点
3. 栅栏
   1. void __threadfence_block() block内栅栏,保证栅栏前被调用线程产生的对共享内存和全局内存的所有写操作对展览后同一块的其他线程是可见的
   2. void __threadfence()网格级内存栅栏,挂起调用的线程,直到全局内存中的所有写操作对相同网格内的所有线程都是可见的
   3. void __threadfence_system()跨系统(包括主机和设备)