Share memory中bank conflict问题

Share memory是片上资源，生命周期是整个block中，它的数据读写十分快，有1个cycle latency。在Share memory中，经常存在bank conflict问题，如果没有bank conflict问题，它的数据读写可以和片上的寄存器（Register）一样快。因此，我们需要尽量减少bank conflicts.
首先，什么是bank？我们以capability 1.x为例，Share memory被等分成同等尺寸大小的存储器模式，即banks，以下图为例：

每一个bank的带宽为32-bit，即4 bytes。连续的32-bit字节可以被分到连续的bank中去。我们以G80为例，它具有16个banks，所以每一个bank=4-byte address % 16，此外，bank的数量和半个warp的thread数量一致，所以share memory对一个warp的请求，分成了两个先后请求来做，不同的半个warp之间没有bank conflicts。如果同时有两个请求在同一个bank中，就会出现bank conflicts，如下图：

对于share memory，有数据读写有快慢两种情况：

快：1)半个warp的所有threads在同一时刻读写share memory的不同bank。2)半个warp中的所有threads在同一时刻以广播的形式读写share memory中的同一个bank。

慢：半个warp的所有threads在同一时刻读写share memory的多个bank(不是全部)，这时必须串行读写。

以G80为例，其share memory具有16个bank，在下面的图示中，只有S为奇数时，才不会存在bank conflicts。


当s=4时，

同样，在存储结构体数据时，也有可能会出现bank conflicts。

例如如下代码存在三路冲突：

struct vector { float x, y, z; };

__shared__ struct vector vectors[32];

struct vector v = vectors[baseIndex + threadIdx.x];

因此，对于share memory，最好的读写方式每个thread读取block中连续的元素。

以如下例程为例：

对于一个二维数组，半个warp的threads读取每一列，这存在了16路bank conflicts。

我们有两种方式可以解决：

1) 在每一行的后面加一个元素，例如：

2)在处理矩阵之前转置矩阵

我们对上一节global memory中的coalescing的例程，作进一步优化，分配额外的一列：

代码如下：

__global__ void transpose_exp(float *odata, float *idata, int width, int height){

__shared__ float block[BLOCK_DIM][BLOCK_DIM+1];

unsigned int xIndex = blockIdx.x * BLOCK_DIM + threadIdx.x;

unsigned int yIndex = blockIdx.y * BLOCK_DIM + threadIdx.y;

if( (xIndex < width)&&(yIndex < height) {

unsigned int index_in = xIndex + yIndex * width;

block[threadIdx.y][threadIdx.x] = idata[index_in];

}

__syncthreads();

xIndex = blockIdx.y * BLOCK_DIM + threadIdx.x;

yIndex = blockIdx.x * BLOCK_DIM + threadIdx.y;

if( (xIndex < height)&&(yIndex < width) ){

unsigned int index_out = yIndex * height + xIndex;

odata[index_out] = block[threadIdx.x][threadIdx.y];

}

}

实验结果为：

 128x128: 0.011ms vs. 0.022ms (2.0X speedup)

 512x512: 0.07ms vs. 0.33ms (4.5X speedup)

 1024x1024: 0.30ms vs. 1.92ms (6.4X speedup)

 1024x2048: 0.79ms vs. 6.6ms (8.4X speedup)

我们再来考虑矩阵相乘时，会不会出现bank conflicts问题，判断技巧在于固定k，考虑同一时刻thread的读些问题

1)当在读一个Ms时，在同一时刻，半个warp的16个threads以广播的形式读取同一个bank

for (int k = 0; k < 16; ++k) //ty is constant over a half-warp
Csub += Ms[ty][k] * Ns[k][tx];//here k is fixed over a half-warp

2)当在读一个Ns时，在同一时刻，半个warp的16个threads连续读取不同的bank，stride=1.

for (int k = 0; k < 16; ++k) //tx is constant over a half-warp
Csub += Ms[ty][k] * Ns[k][tx];

所以综上所述，矩阵相乘的例子并没有bank conflicts问题。

需要注意的是，我们上述所说的均为capability 1.x的情况，在2.x 3.x中，每个share memory有32个bank，具体还是要看技术手册的。