共享存储器的物理模型以及存储数据的方式如下图1所示,每个共享存储器被分为32个bank,可以通过线程块里的线程访问共享存储器,最理想的情况下一个warp(32个线程)一次访问32个bank。图1展示了每个元素4个字节,共1024个元素的一维数据存储于2维共享存储器存储存储的一种结构。为了利于理解和说明,假设图一上面的数据由0依次增加到1023,这些数据按照行为主的顺序存储到共享存储器中,由于共享存储器有32个bank,每个bank可存储32位数据,因而可排列图1下方的结构。
图1
2维静态共享存储器的声明如下:
_ _ shared _ _ int tile[N][N];
其中第一个N表示的为y轴上维度为N,第二个N表示的为x轴上维度为N
二维线程获取共享存储器上数据的方法如下:
tile[threadIdx.y][threadIdx.x]
tile[threadIdx.x][threadIdx.y]
根据之前介绍的共享存储器的硬件模型和最理想的访问方式,第一种二维线程的访问方式更好,不会照成bank冲突,而第二种二维线程访问方式会照成32路bank冲突。
参考资料:Professional CUDA C Programming 及其chapter05 checkSmemSquare.cu代码