【CUDA学习】共享存储器

下面简单介绍一些cuda中的共享存储器和全局存储器　

共享存储器，shared memory，能够被同一块中的全部线程访问的可读写存储器，生存期是块的生命期。

Tesla的每一个SM拥有16KB共享存储器。

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

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

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

共享存储器被组织为16个bank，每一个bank拥有32bit的宽度。

无bank conflict时，一个half-warp内的线程能够在一个内核周期中并行访问

对同一bank的同时访问致使bank conflict 　　只能顺序处理　访存效率下降

若是half-warp的线程访问同一地址时，会产生一次广播，不会产生bank conflict

__shared__ float shared[256];

float foo = shared[threadIdx.x];

没有访问冲突

__shared__ float shared[256];

float foo = shared[threadIdx.x * 2];

产生2路访问冲突

__shared__ float shared[256];

float foo = shared[threadIdx.x*8];

产生8路访问冲突