CUDA之Shared memory bank conflicts详解

目前 CUDA 装置中，每个 multiprocessor 有 16KB 的 shared memory。Shared memory 分成16 个 bank。如果同时每个 thread 是存取不同的 bank，就不会产生任何问题，存取 shared memory 的速度和存取寄存器相同。不过，如果同时有两个（或更多个） threads 存取同一个bank 的数据，就会发生 bank conflict，这些 threads 就必须照顺序去存取，而无法同时存取shared memory 了。

Shared memory 是以 4 bytes 为单位分成 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参照上一篇博客），因此分属于不同的 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 1 和 thread 5 也是同样，这样就会造成 bank conflict。 在这个例子中，一个 half warp 的 16 个 threads 会有四个 threads 存取同一个 bank，因此存取 share memory 的速度会变成原来的 1/4。

一个重要的例外是，当多个 thread 存取到同一个 shared memory 的地址时，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 了。

http://blog.csdn.net/sunmc1204953974/article/details/51078818

bank conflicts 如何产生的？

对GPU来说，local memory是由banks组成的，每个bank是32bit，可视化图如下。bank是实际存储单元。每个bank在一次访存中，可以被取址一次。并行访问相同的bank，将导致访存冲突（bank conflicts）。

Bank    |      1      |      2      |      3      |...
Address |  0  1  2  3 |  4  5  6  7 |  8  9 10 11 |...
Address | 64 65 66 67 | 68 69 70 71 | 72 73 74 75 |...

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

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

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

那么问题来了，为什么 shared memory 存在 bank  conflict，而 global memory 不存在？因为访问 global memory 的只能是 block，而访问 shared memory 的却是同一个 half-warp 中的任意线程。

引用风辰CUDA入门教程的一段话：

Tesla 的每个 SM 拥有 16KB 共享存储器，用于同一个线程块内的线程间通信。为了使一个 half-warp 内的线程能够在一个内核周期中并行访问，共享存储器被组织成 16 个 bank，每个 bank 拥有 32bit 的宽度，故每个 bank 可保存 256 个整形或单精度浮点数，或者说目前的bank 组织成了 256 行 16 列的矩阵。如果一个 half-warp 中有一部分线程访问属于同一bank 的数据，则会产生 bank conflict，降低访存效率，在冲突最严重的情况下，速度会比全局显存还慢，但是如果 half-warp 的线程访问同一地址的时候，会产生一次广播，其速度反而没有下降。在不发生 bank conflict 时，访问共享存储器的速度与寄存器相同。在不同的块之间，共享存储器是毫不相关的。

里面说的很清楚的一点就是每个bank有1KB的存储空间。

Shared memory 是以 4 bytes 为单位分成 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。

const int tid = threadIdx.x;  

因此，如果程序在存取 shared memory 的时候，使用以下的方式：

int number = data[base + tid];

那就不会有任何 bank conflict，可以达到最高的效率。但是，如果是以下的方式：

int number = data[base + 4 * tid];

那么，thread 0 和 thread 4 就会存取到同一个 bank，thread 1 和 thread 5 也是同 样，这样就会造成 bank conflict。在这个例子中，一个 half warp 的 16 个 threads 会有四个threads 存取同一个 bank，因此存取 share memory 的速度会变成原来的 1/4。

CUDA编程中,一个half-warp（16个threads）访问连续的32bit地址,不会有bank conflicts。 一个重要的例外是，当多个 thread 存取到同一个 shared memory 的地址时，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个bank中。一个block要访问shared memory，只要能够保证以其中相邻的16个线程一组访问thread，每个线程与bank是一一对应就不会产生bank conflict。否则会产生bank conflict，访存时间成倍增加，增加的倍数由一个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]; 

则不会产生块访问冲突。

http://blog.163.com/wujiaxing009@126/blog/static/71988399201701610022165/