#01 前情回顾
前面的文章中，我们简单分析了shared memory Bank Conflict的基础概念，本文我们来对如何避免Bank冲突进行分析。
需要注意的是，不同warp中的线程之间不存在Bank冲突，后续的所有分析均是针对一个warp内的线程。
#02 地址映射
要解决bank冲突，首先我们要了解一下共享内存的地址映射方式。在共享内存中，连续的32-bits字被分配到连续的32个bank中，这就像电影院的座位一样：一列的座位就相当于一个bank，所以每行有32个座位，在每个座位上可以“坐”一个32-bits的数据(或者多个小于32-bits的数据，如4个char型的数据，2个short型的数据)；而正常情况下，我们是按照先坐完一行再坐下一行的顺序来坐座位的，在shared memory中地址映射的方式也是这样的。下图中内存地址是按照箭头的方向依次映射的：
上图中数字为bank编号。这样的话，如果申请一个共享内存数组(假设是int类型)的话，那么你的每个元素所对应的bank编号就是地址偏移量(也就是数组下标)对32取余所得的结果，比如大小为1024的一维数组myShMem：
myShMem[4]: 对应的bank id为#4 (相应的行偏移量为0)
myShMem[31]: 对应的bank id为#31 (相应的行偏移量为0)
myShMem[50]: 对应的bank id为#18 (相应的行偏移量为1)
myShMem[128]: 对应的bank id为#0 (相应的行偏移量为4)
myShMem[178]: 对应的bank id为#18 (相应的行偏移量为5)
#03 典型访问方式
步长(stride)为1的访问：由于每个warp中的线程ID与每个bank的ID一一对应，因此不会产生bank冲突。
下面这种访问虽然是交叉的访问，每个线程并没有与bank一一对应，但每个线程都会对应一个唯一的bank，所以也不会产生bank冲突。
下面这种虽然也是线性的访问bank，但这种访问方式与第一种的区别在于访问的步长(stride)变为2，这就造成了线程0与线程28都访问到了bank 0，线程1与线程29都访问到了bank 2...，于是就造成了2路的bank冲突。
下面这种访问造成了8路的bank冲突，
这里我们需要注意，下面这两种情况是两种特殊情况：
上图中，所有的线程都访问了同一个bank，貌似产生了32路的bank冲突，但是由于广播(broadcast)机制(当一个warp中的所有线程访问一个bank中的同一个字(word)地址时，就会向所有的线程广播这个字(word))，这种情况并不会发生bank冲突。
同样，这种访问方式也不会产生bank冲突：
这就是所谓的多播机制(multicast)——当一个warp中的几个线程访问同一个bank中的相同字地址时，会将该字广播给这些线程。
#04 数据类型与bank冲突
由上面的分析，可以看出，当每个线程访问一个32-bits大小的数据类型的数据(如int，float)时，不会发生bank冲突。
extern __shared__ int shrd[];
foo = shrd[baseIndex + threadIdx.x]
但是如果每个线程访问一个字节(8-bits)的数据时，会不会发生bank冲突呢？很明显这种情况会发生bank冲突的，因为四个线程访问了同一个bank，造成了四路bank冲突。同理，如果是short类型（16-bits）也会发生bank冲突，会产生两路的bank冲突，下面是这种情况的两个例子：
extern __shared__ char shrd[];
foo = shrd[baseIndex + threadIdx.x];
extern __shared__ short shrd[];
foo = shrd[baseIndex + threadIdx.x];
#05 访问步长与bank冲突
我们通常这样来访问数组：每个线程根据线程编号tid与s的乘积来访问数组的32-bits字(word):
extern __shared__ float shared[];
float data = shared[baseIndex + s * tid];
如果按照上面的方式，那么当s*n是bank的数量(即32)的整数倍时或者说n是32/d的整数倍(d是32和s的最大公约数)时，线程tid和线程tid+n会访问相同的bank。我们不难知道如果tid与tid+n位于同一个warp时，就会发生bank冲突，相反则不会。
仔细思考你会发现，只有warp的大小(即32)小于等于32/d时，才不会有bank冲突，而只有当d等于1时才能满足这个条件。要想让32和s的最大公约数d为1，s必须为奇数。于是，这里有一个显而易见的结论：当访问步长s为奇数时，就不会发生bank冲突。
#06 bank冲突的示例
下面我们以并行计算中的经典的归约算法为例来做一个简单的练习。
假设有一个大小为2048的向量，我们想用归约算法对该向量求和。于是我们申请了一个大小为1024的线程块，并声明了一个大小为2048的共享内存数组，并将数据从全局内存拷贝到了该共享内存数组。
我们可以有以下两种方式实现归约算法：
不连续的方式：
连续的方式：
下面我们用具体的代码来实现上述两种方法
// 非连续的归约求和
__global__ void BC_addKernel(const int *a, int *r)
{
__shared__ int cache[ThreadsPerBlock];
int tid = blockIdx.x * blockDim.x + threadIdx.x;
int cacheIndex = threadIdx.x;
// copy data to shared memory from global memory
cache[cacheIndex] = a[tid];
__syncthreads();
// add these data using reduce
for (int i = 1; i < blockDim.x; i *= 2)
{
int index = 2 * i * cacheIndex;
if (index < blockDim.x)
{
cache[index] += cache[index + i];
}
__syncthreads();
}
// copy the result of reduce to global memory
if (cacheIndex == 0)
r[blockIdx.x] = cache[cacheIndex];
}
上述代码实现的是非连续的归约求和，从int index = 2 * i * cacheIndex和cache[index] += cache[index + i];两条语句，我们可以很容易判断这种实现方式会产生bank冲突。当i=1时，步长s=2xi=2，会产生两路的bank冲突；当i=2时，步长s=2xi=4，会产生四路的bank冲突...当i=n时，步长s=2xn=2n。可以看出每一次步长都是偶数，因此这种方式会产生严重的bank冲突。
连续的归约求和
// 连续的归约求和
__global__ void NBC_addKernel2(const int *a, int *r)
{
__shared__ int cache[ThreadsPerBlock];
int tid = blockIdx.x * blockDim.x + threadIdx.x;
int cacheIndex = threadIdx.x;
// copy data to shared memory from global memory
cache[cacheIndex] = a[tid];
__syncthreads();
// add these data using reduce
for (int i = blockDim.x / 2; i > 0; i /= 2)
{
if (cacheIndex < i)
{
cache[cacheIndex] += cache[cacheIndex + i];
}
__syncthreads();
}
// copy the result of reduce to global memory
if (cacheIndex == 0)
r[blockIdx.x] = cache[cacheIndex];
}
