下面简单介绍一些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路访问冲突
