cuda合并访问的要求_CUDA 共享内存的bank conflict 以及合并存储器访问（Coalesced memory accesses）...

上面两个概念不要搞混了，两个不同的概念

大部分转自http://www.cnblogs.com/waytofall/archive/2013/02/19/2916996.html

http://www.cnblogs.com/dwdxdy/p/3215187.html

个人感觉shared memory

可能是CUDA优化编程中最经常考虑的东西了。

在编程过程中，有静态的shared memory 动态的shared memory

静态的shared memory 在程序中定义　　　__shared__ type shared[SIZE];

动态的shared memory 通过内核函数的每三个参数设置大小 extern __shared__ type

shared[];

为什么 shared memory 存在 bank conflict，而 global memory 不存在？因为访问 global memory 的只能是 block，而访问 shared memory 的却是同一个 half-warp 中的任意线程。

Tesla 的每个 SM 拥有 16KB 共享存储器，用于同一个线程块内的线程间通信。为了使一个 half-warp 内的线程能够在一个内核周期中并行访问，共享存储器被组织成 16 个 bank，每个 bank 拥有 32bit 的宽度，故每个 bank 可保存 256 个整形或单精度浮点数，或者说目前的bank 组织成了 256 行 16 列的矩阵。如果一个 half-warp 中有一部分线程访问属于同一bank 的数据，则会产生 bank

conflict，降低访存效率，在冲突最严重的情况下，速度会比全局显存还慢，但是如果 half-warp 的线程访问同一地址的时候&#x