Share Memory & Bank Conflict

2024-03-08 13:49:04

共享存储器

share memory是MP拥有的资源，因为它是一个片上存储器，所以访问共享存储器比本地存储器和全局存储器要快得多。实际上共享存储器的延迟大约比没有缓存的全局存储器低100倍(假设线程之间没有bank冲突)。共享存储器被分配给每个线程块，所以块内的线程可以访问同一个共享存储器。线程可以访问共享内存中由同一线程块中的其他线程从全局内存加载的数据

线程同步

当在线程之间共享数据时，我们需要小心以避免竞态条件(race conditions)。因为线程块中的线程之间虽然逻辑上是并行的，但是物理上并不是同时执行的。让我们假设线程A和线程B分别从全局存储器中加载了一个数据并且将它存到了共享存储器。然后，线程A想要从共享存储器中读取B的数据，反之亦然。我们还要假设线程A和B位于不同的warp。如果在A尝试读取B的数据时，B还未写入，这样就会导致未定义的行为和错误的结果。

为了保证在并行线程协作时得到正确的结果，我们必须对线程进行同步。CUDA提供了一个简单的栅栏同步原语，__syncthreads()。每个线程只能在块中所有的线程执行完__syncthreads()函数后，才能继续执行__syncthreads()的语句。因此我们可以在向共享存储器存数据后以及在向共享存储器加载数据前调用__syncthreads()，这样就避免了上面所描述的竞态条件(race conditions)。我们必须要牢记 __syncthreads()被用在分支代码块中是未定义的行为，很可能会导致死锁——线程块中所有的线程必须在同一点调用__syncthreads()

例子：数组逆序

在设备代码中声明共享内存要使用__shared__变量声明说明符。两种方式申请共享内存，不同之处在于共享内存数组的声明以及核函数的调用。：

1. 静态共享内存， 大小在编译时可确定
2. 动态共享内存， 大小在运行时确定

静态共享内存

    __global__ void staticReverse(int *d, int n)
    {
      __shared__ int s[64];  //静态共享内存数组声明，数组长度在编译时就确定
      int t = threadIdx.x;
      int tr = n-t-1;
      s[t] = d[t];
      __syncthreads();
      d[t] = s[tr];
    }
    
    int main(){
        ...
        staticReverse <<< 1, n >>> (d_d, n);
        ...
    }

静态共享内存数组声明，数组长度在编译时就确定。

动态共享内存

__global__ void dynamicReverse(int *d, int n)
{
  extern __shared__ int s[]; //静态共享内存数组声明，数组长度在运行译时确定
  int t = threadIdx.x;
  int tr = n-t-1;
  s[t] = d[t];
  __syncthreads();
  d[t] = s[tr];
}

int main(){
    ...
    dynamicReverse <<<1 , n ,n * sizeof(int) >>> (d_d, n);
    ...
}

静态共享内存数组声明时，需要加 extern 关键字。
当核函数被启动时，数组大小从第三个执行配置参数被隐式地确定。

在一个核函数中动态地申请多个数组

如果你想在一个核函数中动态地申请多个数组时该怎么办呢？你必须在首先申请一个单独的未指定大小的extern数组，然后使用指针将它分为多个数组，如下所示：

extern __shared__ int s[];
int *integerData = s;                        // nI ints
float *floatData = (float*)&integerData[nI]; // nF floats
char *charData = (char*)&floatData[nF];      // nC chars

这样的话，你需要在核函数中这样指定共享内存的大小：

myKernel<<<gridSize, blockSize, nI*sizeof(int)+nF*sizeof(float)+nC*sizeof(char)>>>(...);

bank conflict

参考

对不同bank的访问可同时进行 ：为了获得较高的内存带宽，共享存储器被划分为多个大小相等的存储器模块，称为bank，可以被同时访问。因此任何跨越b个不同bank的n个地址的读写操作可以被同时进行，这样就大大提高了整体带宽 ——可达到单独一个bank带宽的b倍。

同一warp中的所有线程访问同一bank可同时进行（广播） ：然而，如果多个线程请求的地址映射到相同的内存bank，那么访问就会被顺序执行。硬件会把冲突的内存请求分为尽可能多的单独的没有冲突的请求，这样就会减少一定的带宽，减少的因子与冲突的内存请求个数相等。当然，也有例外的情况：当一个warp中的所有线程访问同一个共享内存地址时，就会产生一次广播。计算能力为2.0及以上的设备还可以多次广播共享内存访问，这意味着一个warp中任意数量的线程对于同一位置的多次访问也可以同时进行。

为了尽量减少bank冲突，理解共享内存地址如何映射到bank是非常重要的。计算能力5.0的设备，共享内存的bank是这样组织的：连续的32-bits字被分配到连续的bank中，每个bank的带宽是每个时钟周期32bits。；对于计算能力3.0的设备，每个bank的带宽是每个时钟周期64bits；对于计算能力2.0的设备每个bank的带宽是每两个时钟周期32bits

注释：

对于计算能力1.x的设备，warp的大小是32而bank的数量是16。一个warp中线程对共享内存的请求被划分为两次请求：一个请求是前半个warp的另一个请求时后半个warp的。注意如果每个bank中只有一个内存地址是被半个warp中的线程访问的话，是不会有bank冲突的。
对于计算能力为2.x的设备，warp的大小是32而bank的数量也是32。一个warp中线程对共享内存的请求不会像计算能力1.x的设备那样被划分开，这就意味着同一个warp中的前半个warp中的线程与后半个warp中的线程会有可能产生bank冲突的。
计算能力为3.x的设备的bank大小是可以配置的，我们可以通过函数cudaDeviceSetSharedMemConfig()来设置，要么设置为4字节(默认为cudaSharedMemBankSizeFourByte)，要么设置为8字节(cudaSharedMemBankSizeEightByte)。当访问双精度的数据时，将bank大小设置为8字节可以帮助避免bank冲突。

当一个warp中的不同线程访问一个bank中的不同的字地址时，就会发生bank冲突。

如果没有bank冲突的话，共享内存的访存速度将会非常的快，大约比全局内存的访问延迟低100多倍，但是速度没有寄存器快。然而，如果在使用共享内存时发生了bank冲突的话，性能将会降低很多很多。在最坏的情况下，即一个warp中的所有线程访问了相同bank的32个不同字地址的话，那么这32个访问操作将会全部被序列化，大大降低了内存带宽。

注意：不同warp中的线程之间不存在什么bank冲突。

要解决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)

典型的bank访问方式

发生bank冲突的情况：

下面这这种访问方式是典型的线性访问方式(访问步长(stride)为1)，由于每个warp中的线程ID与每个bank的ID一一对应，因此不会产生bank冲突。

下面这种访问虽然是交叉的访问，每个线程并没有与bank一一对应，但每个线程都会对应一个唯一的bank，所以也不会产生bank冲突。

下面这种虽然也是线性的访问bank，但这种访问方式与第一种的区别在于访问的步长(stride)变为2，这就造成了线程0与线程28都访问到了bank 0，线程1与线程29都访问到了bank 2...，于是就造成了2路的bank冲突。我在后面会对以不同的步长(stride)访问bank的情况做进一步讨论。

下面这种访问造成了8路的bank冲突，

没有bank冲突的情况：

这里我们需要注意，下面这两种情况是两种特殊情况：

上图中，所有的线程都访问了同一个bank，貌似产生了32路的bank冲突，但是由于广播(broadcast)机制(当一个warp中的所有线程访问一个bank中的同一个字(word)地址时，就会向所有的线程广播这个字(word))，这种情况并不会发生bank冲突。

同样，这种访问方式也不会产生bank冲突：

这就是所谓的多播机制(multicast)——当一个warp中的几个线程访问同一个bank中的相同字地址时，会将该字广播给这些线程。

NOTE: 这里的多播机制(multicast)只适用于计算能力2.0及以上的设备。

数据类型与bank冲突

我们都知道，当每个线程访问一个32-bits大小的数据类型的数据(如int，float)时，不会发生bank冲突。

extern __shared__ int shrd[];
foo = shrd[baseIndex + threadIdx.x]

但是如果每个线程访问一个字节(8-bits)的数据时，会不会发生bank冲突呢？其实这种情况是不会发生bank冲突的。当同一个字(word)中的不同字节被访问时，也不会发生bank冲突，下面是这种情况的两个例子：

extern __shared__ char shrd[];
foo = shrd[baseIndex + threadIdx.x];

extern __shared__ short shrd[];
foo = shrd[baseIndex + threadIdx.x];

访问步长与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冲突。

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];
}

由于每个线程的ID与操作的数据编号一一对应，因此上述的代码很明显不会产生bank冲突。

码农公寓