CUDA 共享内存

GPU设备中的共享内存比全局内存访问效率高，一种借助共享内存的通用办法是，数据分块之后，将一个线程块中需要从全局内存中读取的数据全部读到共享内存中，然后线程块中的所有线程在执行时，只需要读取共享内存即可。

共享内存中所谓的共享是对线程块而言，也就是说同一个线程块中的所有线程共享这块内存。为了使得一个wrap中的线程可以在同一个时刻并行访问共享内存，共享存储器被组织成了32个bank，每个bank有4字节。

共享内存访问有bank 冲突问题，一个wrap中存在多个线程同时访问同一个bank的数据，则称为bank冲突，bank冲突非常影响性能，在冲突最严重的情况下，访问速度比全局内存还要慢。但是如果一个wrap中多有线程访问同一个bank数据，不会产生bank冲突，而是会产生一次广播，访问速度不会下降。在不发生bank冲突时，访问共享存储器的速度与寄存器相当。当然在不同块之间，共享存储器是毫不相关的。

假设定义以下数据：

__shared__ float data[128];

那么，data[0]对应bank0，data[1]对应bank1，data[2]对应bank2，........data[31]对应bank31，而data[32]又回到bank0，data[33]回到bank1.

因此，如果程序以如下方式访问：

foat f = data[base+tid]；则不会产生bank 冲突，因为各自线程访问不同的bank。

但是如果以下方式访问:

float f = data[base + 4*tid]；thread0和thread8 就会访问bank0，thre