CUDA中Bank conflict冲突
转自:http://blog.csdn.net/smsmn/article/details/6336060
其实这两天一直不知道什么叫bank conflict冲突,这两天因为要看那个矩阵转置优化的问题,里面有讲到这些问题,但是没办法,为了要看懂那个bank conflict冲突,我不得不去找资料,说句实话我现在不是完全弄明白,但是应该说有点眉目了,现在我就把网上找的整理一下,放在这边,等哪天完全弄明白了我就在修改里面的错误。
Tesla 的每个 SM 拥有 16KB 共享存储器,用于同一个线程块内的线程间通信。这里先假设一个SM里头有一个block。则每个block为1KB,即1024*8 bits。为了使一个 half-warp 内的线程能够在一个内核周期中并行访问,共享存储器被组织成 16 个 bank,每个 bank 拥有 32bit 的宽度,1024*8 bits = 32 * 256bits,32bits=4bytes=1个int数据的大小,故每个 bank 可保存 256 个整形或单精度浮点数,或者说目前的bank 组织成了 256 行 16 列的矩阵(这里16应该是32才对?)。如果一个 half-warp 中有一部分线程访问属于同一bank 的数据,则会产生 bank conflict,降低访存效率,在冲突最严重的情况下,速度会比全局显存还慢,但是如果 half-warp 的线程访问同一地址的时候,会产生一次广播,其速度反而没有下降。在不发生 bank conflict 时,访问共享存储器的速度与寄存器相同。在不同的块之间,共享存储器是毫不相关的。 ------风辰的 CUDA 入门教程
里面说的很清楚就是每个bank有1KB的存储空间。
Shared memory 是以 4 bytes (32bits, int 是32bits, 4byte)为单位分成 banks。因此,假设以下的数据:
__shared__ int data[128];
那么,data[0] 是 bank 0、data[1] 是 bank 1、data[2] 是 bank 2、…、data[15] 是bank 15,而 data[16] 又回到 bank 0。由于 warp 在执行时是以 half-warp 的方式执行,因此分属于不同的 half warp 的 threads,不会造成 bank conflict。
因此,如果程序在存取 shared memory 的时候,使用以下的方式:
int number = data[base + tid];
那就不会有任何 bank conflict,可以达到最高的效率。但是,如果是以下的方式:
int number = data[base + 4 * tid];
那么,thread 0 和 thread 4 就会存取到同一个 bank(意思是说,thread 0对应的是data[base],thread 4对应的是data[base+16],就是对应同一个bank了,但是是不同行),thread 1 和 thread 5 也是同 样,这样就会造成 bank conflict。在这个例子中,一个 half warp 的 16 个 threads 会有四个threads 存取同一个 bank,因此存取 share memory 的速度会变成原来的 1/4。
一个重要的例外是,当多个 thread 存取到同一个 shared memory 的地址时(即同一个bank,shared memory 存储器是被划分为16个小单元,与half-warp长度相同,称为bank,每个bank可以提供自己的地址服务。),shared memory 可以将这个地址的 32 bits 数据「广播」到所有读取的 threads,因此不会造成 bank conflict。例如:
int number = data[3];
这样不会造成 bank conflict,因为所有的 thread 都读取同一个地址的数据。
很多时候 shared memory 的 bank conflict 可以透过修改数据存放的方式来解决。例如,以下的程序:
data[tid] = global_data[tid];
...
int number = data[16 * tid];
会造成严重的 bank conflict,为了避免这个问题,可以把数据的排列方式稍加修改,把存取方式改成:
int row = tid / 16;
int column = tid % 16;
data[row * 17 + column] = global_data[tid];
...
int number = data[17 * tid];
这样就不会造成 bank conflict 了。
简单的说,矩阵中的数据是按照bank存储的,第i个数据存储在第i%16个bank中。一个block要访问shared memory,只要能够保证以其中相邻的16个线程一组访问thread,每个线程与bank是一一对应就不会产生bank conflict。否则会产生bankconflict,访存时间成倍增加,增加的倍数由一个bank最多被多少个thread同时访问决定。有一种极端情况,就是所有的16个thread同时访问同一bank时反而只需要一个访问周期,此时产生了一次广播。
下面有一些小技巧可以避免bank conflict 或者提高global存储器的访问速度
1. 尽量按行操作,需要按列操作时可以先对矩阵进行转置
2. 划分子问题时,使每个block处理的问题宽度恰好为16的整数倍,使得访存可以按照 s_data[tid]=i_data[tid]的形式进行
3. 使用对齐的数据格式,尽量使用nvidia定义的格式如float3,int2等,这些格式本身已经对齐。
4. 当要处理的矩阵宽度不是16的整数倍时,将其补为16的整数倍,或者用malloctopitch而不是malloc。
5. 利用广播,例如s_odata[tid] = tid%16 < 8 ? s_idata[tid] : s_idata[15];会产生8路的块访问冲突而用:s_odata[tid]=s_idata[15];s_odata[tid]= tid%16 < 8 ? s_idata[tid] : s_data[tid]; 则不会产生块访问冲突